cuda注意

1，不能在主机代码中对cudamalloc()返回的指针进行解应用，主机代码可以将这个指针作为参数传递，对其进行算术运算，甚至可以将其转换为另一种不同的类型，但是，不可使用这个指针来读取或者写入内存；

2，主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存；

3，memPitch:在内存复制中最大的修正量，单位是字节，在cudaDeviceProp中的字段。作用；

4，GPU中通常使用的分配内存的方式使用cudamalloc，但是对于二维和三维的矩阵而言，使用cudaMalloc来分配内存并不能的得到最后的性能，原因是：对于二维和三维的内存，对齐是一个很重要的性质，cudaMallocPitch或者cudaMalloc3D这两个函数能够保证分配的怒才能是合理对齐的，满足物理上的内存访问。可以确保对行访问时具有最优的效率。除此之外，应当使用cudaMemcpy2D和cudaMemcpy3D来实现，

2）cudaMallocPitch()的使用：float *dc_mat = (float*)malloc(sizeof(float)*M*N);  size_t pitch; cudaMallocPitch(&d_mat,&pitch,M*sizeof(float),N);之后调用kernel函数，将数据拷贝出来；cudaMemecpy2D(dc_mat,M*sizeof(float),d_mat,pitch,M*sizeof(float),N,cudaMemcpyToHost);这里的pitch是一行内存的大小，即M*sizeof(float)

5,根据设备的属性可以选择使用最多处理器的ＧＰＵ来执行代码，

6，GPU的计算功能在1.3或者更高的显卡才能支持双精度的数学计算

7，在设备查询到时候，一种是直接通过迭代的方式举个设备获得设备信息，另一种是：cudaDeviceProp prop;int dev   cudaGetDevice(&dev);  之后给prop对应的字段赋值：

memset(&prop,0,sizeof(cudaDeviceProp));prop.major = 1;prop.minor = 3;   之后通过cudaChooseDevice(&dev,&prop)，返回匹配设备的编号，之后就选择对应的设备：cudaSetDevice(dev);

8,在核函数中分配的块数，<<<N,1>>>在GPU上是并行进行的，GPU将运行的核函数的N 个副本，在代码中通过blockIdx.x.来说明是哪个线程块在执行。并行线程块的集合称为一个线程格。线程块的每一维最大不能超过65535，这是一种硬件限制。超过了这个范围，就会导致程序运行失败。

9，dim3 grid(DIM ,DIM )用于指定启动的指定的线程块的数量，

10，共享内存与同步：共享内存：GPU上启动的每一个线程块，都将创建该变量的一个副本，线程块中的每个线程都共享这块内存，但是线程不能修改其他线程块中的变量副本，

这样，一个线程块中的多个线程可以进行通信和协作。共享内存留在GPU上，延迟要第很多，

11，规约算法，树状加法（在深入浅出cuda中）；

12，每个multiprocessor对应一个block，每个multiprocessor有8个stream processor,；要计算分配的块数，1）先得到GPU有多少个流多处理器，2）每个流多处理器有多少的寄存器，3）每个线程需要多少寄存器，得到每个块中分配多少线程；每个线程块中的线程是32的倍数，符合warp为32的要求，

13，3.7 寄存器
寄存器是NVIDIA GPU存储器空间中，最快的硬件，读取它几乎不用耗费时间。因此合理的使用它是至关重要的。

3.7.1 寄存器溢出和本地存储器
由于ＧＰＵ上寄存器的数目是有限的，在计算能力1.0,1.1的设备上，其数目是8 K，在计算能力1.2,1.3的设备上，其值是16 K。如果我们使用的寄存器数目超过了系统的最大数目，编译器便会将其转入设备存储器上的一个区域内，这个区域称为本地存储器。本地存储器的速度和全局存储器一样。
使用2.3及以前的toolkit,可以使用-keep指令输出中间文件，然后直接查看cubin文件就可以知道内核使用的寄存器数目，同时也能看到共享存储器、常量存储器和本地存储器的使用量。如果使用的是3.0的toolkit，此时cubin文件已经不是文本格式，因此不能直接查看，此时可以在用nvcc 编译的时候，使用—ptxas-options=-v选项，编译器会报告存储器使用信息。

3.8 执行配置和占用率
使用<<<grid, block>>>语法指定执行线程配置的时候，grid 和block大小也影响程序的效率。一般而言，grid要大于多处理器的数目，这样才能让多处理器不至于空闲，但是这样也会导致一些问题，比如负载均衡，如果grid大小不能比sm 数目整除的话，就会有ＳＭ计算的时候，另外一些ＳＭ空闲，如果grid远大于ｓｍ数目的话，可忽略，但是如果ＳＭ数目与ｇｒｉｄ大小相差不大的话，性能损耗就很可观了。一般而言，ｇｒｉｄ大小至少是ｓｍ数目的三倍，如果有块内同步的话，ｇｒｉｄ要data_len大于ｓｍ的四倍以上。一般而言，ｂｌｏｃｋ大小要是束大小的四倍以上，此时基本上可隐藏访存延迟。如果数据量比较小的话，ｇｒｉｄ大小和ｂｌｏｃｋ大小可能相互牵制，此时要综合考虑。在数据量比较大的时候，只要考虑ｂｌｏｃｋ大小就行了。
另外，CUDA还有一个占用率问题，所谓占用率就是ＳＭ上活跃块数目和最大允许数目的比例。使用ｃｕｄａ　ｖｉｓｕａｌ　ｐｒｏｆｉｌｅｒ的时候会有这个值。这个值的影响因素有内核内寄存器、共享存储器的用量。由于CUDA线程在切换的时候并不转储线程的状态，这是ＣＵＤＡ线程极轻的原因，但是这也使得为了保证线程能够切换，内核内使用的资源最多只能是最大资源的一半，此时占用率是0.25，一般而言占用率要在0.5左右就可以了。

14，__shared__ int data[] :共享内存是以4bytes为单位分成banks，那么：data[0]位data[0]依次类推

15，常量内存用于保存在内核函数执行期间不会发生变化的数据。在某些情况下，用常量内存来替换全局内存能有效的减少内存带宽。、

16，告诉CPU在某个事件上同步，这个事件是cudaEventSynchronize();

17，案例：cudaEvent_t start,stop;  创建事件
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);  记录开始时间

执行的事件

      cudaEventRecord(stop,0);
cudaEventSynchronize(stop);  确保所有的GPU事件都执行完了
float elapsedTime;
HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,start,stop)); 记事件的时间。
printf("the time of before mid_data is %f\n",elapsedTime);
HANDLE_ERROR(cudaEventDestroy(start));销毁事件。防止内存泄露
HANDLE_ERROR(cudaEventDestroy(stop));

18，cuda写循环慢：GPU中有成千上万的线程的吞吐量可以提升速度，但是其中每个线程的运算功能要比cpu慢。CPU擅长处理循环。，

19，常量内存提高性能：对常量内存的单次读取操作可以广播到其他的“邻近”线程，这将节约15次读取操作

2）常量内存的数据将缓存起来，因此对于相同地址的连续操作将不会产生额外的内存通信量。

在函数外定义常量：__constant__ int num[size];     在主函数中使用cudaMemcpyToSymbol(num,src,size);之后在内车函数中使用就可以了。不用在设备和host之间传播。

20，纹理内存：在特定的访问模式中，纹理内存同样能够提升性能并减少内存流量。

2)纹理存储器的声明的大小比常数存储器要大的多。在kernel中访问纹理存储器的操作称为纹理拾取，将显卡中的数据与纹理参照系关联的操作，称为将数据与纹理绑定。

显存中可以绑定到纹理的数据有两种：分为普通的线性存储器和cuda数组。

线性存储器只能与1，2维的纹理绑定，采用整型纹理拾取坐标，坐标值与数据在存储器中的位置相同。

CUDA数组可以与一维和二维和三维纹理绑定，纹理拾取坐标为归一化或者非归一化的浮点型，并且支持许多特殊功能。

纹理缓存中的数据可以被重复利用

纹理缓存一次预取拾取坐标对应位置附近的几个像元，可以实现滤波模式、。

21，有文件表名：的那个线程块的数量为GPU中处理器数量(流多处理器)的2倍的时候，将达到最佳性能。

22.__syncthreads():是cuda的内置函数，头文件为device_function.h.保证block块中可以到达的都到达了。最重要的理解是那些可以到达__syncthreads()的线程需要其他可以到达该点的线程，而不是等待块内所有其他线程。

23，