CUDA编程入门笔记

1、线程块（block）是独立执行的，在执行的过程中线程块之间互不干扰，因此它们的执行顺序是随机的

2、同一线程块中的线程可以通过访问共享内存（shared memory）或者通过同步函数__syncthreads()来协调合作。

3、cuda全称：computer union device architecture    统一计算设备架构，因此CUDA并不是 编程语言

从线程层次看：

从内存层次看：

4、CUDA提供的API主要有两种： runtime API和driver API，其中runtime API实际上是driver API的封装，方便程序员编程，而driver API可以操纵更为底层的东西，例如控制CUDA Contexts（一种类似主机进程的概念）以及CUDA Modules（类似主机动态加载库的概念）等更加底层的CUDA模块。

5、

6、二进制代码在CUDA计算设备上具有小版本的向前兼容性，但是在大版本上不具备兼容性

7、在显卡内存中称为global memory

8、cudaError_t cudaMalloc( void** devPtr，size_t count )函数主要作用是在GPU中申请一块内存地址，向设备分配 count 字节的线性存储器，并以*devPtr的形式返回指向所分配存储器的指针。可针对任何类型的变量合理调整所分配的存储器。存储器不会被清除。如果出现错误，cudaMalloc()将返回cudaErrorMemoryAllocation。

9、cudaError_t cudaMemcpy( void* dst，const void* src，size_t count，enum cudaMemcpyKind kind )函数的主要作用是GPU与CPU之间的数据复制，从src指向的存储器区域中将count个字节复制到dst指向的存储器区域，其中kind是cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost或cudaMemcpyDeviceToDevice之一，用于指定复制的方向。存储器区域不可重叠。调用cudaMemcpy()时，如果dst和src指针与复制的方向不匹配，则将导致不确定的行为。

cudaError_t cudaMemcpyAsync( void* dst，const void* src，size_t count，enum cudaMemcpyKind kind，cudaStream_t stream )
cudaMemcpyAsync()是异步的，可选择传入非零流参数，从而将其关联到一个流。它仅对分页锁定的主存储器有效，如果传入指向可分页存储器的指针，那么将返回一个错误。

10、cudaError_t cudaFree (void* devPtr)释放GPU上的存储器，释放devPtr（必须在之前调用cudaMalloc()或cudaMallocPitch()时返回）指向的存储器空间。如果未返回或者之前已经调用过cudaFree(devPtr)，则返回一个错误。如果devPtr为0，则不执行任何操作。如果出现错误，cudaFree()将返回cudaErrorInvalid-DevicePointer。

11、cudaMalloc、cudaMemcpy和cudaFree三个函数，如果其为之前为异步启动，函数可能会返回错误码，例如：cudaMencpy调用前cudaMalloc还未调用过

12、_global__ void kernel(param list){}核函数只能在主机端调用，调用时必须申明执行参数。调用形式如下：Kernel<<<Dg,Db, Ns, S>>>(param list);
<<<>>>运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。<<<>>>运算符对kernel函数完整的执行配置参数形式是<<<Dg, Db, Ns, S>>>
   参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。
参数Db用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x, Db.y, Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。 一个block中共有Db.x*Db.y*Db.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写。参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流之中。

13、texure也是全局存储器，速度比global还要更快，但是为只读，即对应的矩阵的变化不会影响传入的数组。

14、只要同一个 warp 的不同线程会访问到同一个 bank 的不同地址就会发生 bank conflict，除此之外的都不会发生 bank conflict。

15、cudaMallocPitch(void**,int*,widthInBytes,height):这个函数是在线性内存中分配二维数组，因此在使用时，还是用一维的方式使用。注意，width的单位是字节，而height单位是数据类型，而第二个参数的含义下面一段话说明了。说明 向设备分配至少widthInBytes*height字节的线性存储器，并以*devPtr的形式返回指向所分配存储器的指针。该函数可以填充所分配的存储器，以确保在地址从一行更新到另一行时，给定行的对应指针依然满足对齐要求。
c语言申请2维内存时，一般是连续存放的。a[y][x]存放在第y*widthofx*sizeof(元素)+x*sizeof(元素)个字节。但在cuda的global memory访问中，从256字节对齐的地址(addr=0, 256, 512, ...)开始的连续访问是最有效率的。这样，为了提高内存访问的效率，有了cudaMallocPitch函数。cudaMallocPitch函数分配的内存中，数组的每一行的第一个元素的开始地址都保证是对齐的。因为每行有多少个数据是不确定的，widthofx*sizeof(元素)不一定是256的倍数。故此，为保证数组的每一行的第一个元素的开始地址对齐，cudaMallocPitch在分配内存时，每行会多分配一些字节，以保证widthofx*sizeof(元素)+多分配的字节是256的倍数(对齐)。这样，上面的y*widthofx*sizeof(元素)+x*sizeof(元素)来计算a[y][x]的地址就不正确了。而应该是y*[widthofx*sizeof(元素)+多分配的字节]+x*sizeof(元素)。而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。说明：widthInBytes作为输入参数，应该是widthofx*sizeof(元素)；这样的话，复制内容时也要作相应的修改。