# 项目首页

欢迎访问我的 Sphinx 项目！以下是本项目的章节内容。

```{toctree}
:maxdepth: 2
:caption: "目录"

CUDA-C-Best-Practices-Guide/index.md
CUDA-C-Programming-Guide/index.md
```


## CUDA编程基本知识

CUDA Runtime API：函数以`cuda`开头。

CUDA Driver API：函数以`cu`开头。


### 延迟（Latency）

#### 指令延迟

| 数据类型 	| 字长（bits） 	|       操作      	| 延迟（cycles） 	|                                          备注                                          	|
|:--------:	|:------------:	|:---------------:	|:--------------:	|:--------------------------------------------------------------------------------------:	|
| 整数     	|              	| +               	| 4              	|                                                                                        	|
| 整数     	| 32           	| *               	| 16             	|                                                                                        	|
| 整数     	|              	| /,mod           	| 特别昂贵       	| 尽可能避免，或使用位操作替代                                                           	|
|          	|              	| 位操作          	| 4              	|                                                                                        	|
|          	|              	| 比较            	| 4              	|                                                                                        	|
|          	|              	| max,min         	| 4              	|                                                                                        	|
|          	|              	| 类型转换        	| 4              	|                                                                                        	|
| 浮点     	|              	| +,*+            	| 4              	|                                                                                        	|
| 浮点     	|              	| /               	| 36             	|                                                                                        	|
|          	|              	| 倒数            	| 16             	|                                                                                        	|
|          	|              	| 平方根倒数      	| 16             	|                                                                                        	|
| 浮点     	|              	| 平方根          	| 32             	| 实现：平方根倒数+倒数，对0和无穷大能正确得到结果；因此它处理一个warp花费32个时钟周期。 	|
|          	|              	| __log(x)        	| 16             	|                                                                                        	|
|          	|              	| __sin(x)        	| 32             	|                                                                                        	|
|          	|              	| __cos(x)        	| 32             	|                                                                                        	|
|          	|              	| __exp(x)        	| 32             	|                                                                                        	|
|          	|              	| __fdividef(x,y) 	| 20             	| 除法                                                                                   	|
| -        	| -            	| 其他函数        	| 更多时钟周期   	| 因为实现为多个指令的组合                                                               	|
| -        	| -            	| 内存指令        	| 4              	| 包括从共享或全局内存中读写的任何指令，SM使用4个时钟周期来发射warp的一个内存指令        	|
| -        	| -            	| __syncthreads() 	| 4              	| 同步                                                                                   	|
 
有时候，编译器不得不额外插入转换指令，从而引入额外的执行周期：
>对操作数为char或short的函数操作，其操作数通常转换为int；  
>在单精度浮点数计算中，将双精度浮点数常量（不使用任何类型后缀定义）作为输入值；  
>在数学函数的双精度版本，将单精度浮点数变量作为输入参数。

后两种情况可以通过以下方式避免：  
>对单精度浮点数常量，使用f后缀定义，比如3.141592653f  
>对数学函数的单精度版本，也使用f后缀定义，比如sinf().

#### 访存延迟

| 内存类型         	| 延迟    	|
|------------------	|---------	|
| L1/Shared Memory 	| 10-20   	|
| Global Memory    	| 400-600 	|


### 性能指南

《CUDA中文手册》第5章