cuda 学习笔记

1. __global__: Alerts the compiler that a function should be compiled to run on a device instead of the host.

2.The trick is actually in calling the device code from the host code which is handled by the CUDA compiler and runtime.

3.The hardware limits the number of blocks in a single launch to 65535. Similarly, the hardware limits the number of threads per block

  with which we can launch a kernel by the maxThreadPerBlock field of the device properties structure.

4.Cuda C compiler treats variables in shared memory differently than typical variables. It creates a copy of the variable for each block that you launch on the GPU. Every thread in that block shares the memory, but threads cannot see or modify the copy of this variable that is seen within other blocks.
5. Without synchronization, we have created a race condition where the correctness of the execution results depends on the nondeterministic details of the hardware.

6.一种关于__syncthreads() 的无穷等待的情况。关键是认识到：The Cuda Architecture guarantees that no thread will advance to an instruction beyond the __syncthreads() until every thread in the block has executed __syncthreads(). For example:

      if(cacheIndex<i) {

       cache[cacheIndex] += cache[cacheIndex  + i];

      __syncthreads();

}

 如果这样写的话，因为不符合条件的线程不会执行到__syncthreads();而cuda runtime 会一直等待所有线程（包括这种不符合条件的线程）执行到__syncthreads();就会出现无穷等待的情况。

 

7.Reading from constant memory can conserve memory bandwidth when compared to reading the same data from global memory.

 8. The trickiest part of using events arises as a consequence of the fact that some of the calls we make in CUDA C are actually asynchronous. For example, when we launch the kernel in our ray tracer, the GPU begins executing our code, but the CPU continues executing the next line of our program before the GPU finishes.v 

 

9.为什么Kernel函数没有返回值？ 在主机提交Kernel一段时间后，Kernel才开始在GPU上实际投入运行。这种异步调用机制导致Kernel无法返回函数值。

 

10.Cuda 编程的三个基本法则：

    1）将数据放入并始终存储于GPGPU

    2) 交给GPGPU足够更多的任务

    3）注重GPGPU上的数据重用，以避免带宽限制

 

11.可用于CUDA计算的GPGPU包含片上与片外两大类存储器。流多处理器片上存储器是速度最快，可扩展性最佳的备受青睐的存储器。这些存储器容量有限，通常仅在KB级别。板载的全局内存是一个共享的存储系统，可以被GPU上的所有流多处理器访问。该内存容量可以达到GB级别，是目前最大，使用最普遍，但也是GPU上最慢的存储器。

12.一个定义为volatile的变量是说这个变量可能会被意想不到地改变，这样，编译器就不会去假设这个变量的值了。精确的说就是，优化器在用到这个变量的时候必须每次都小心的重新读取这个变量的值，而不是使用保存在寄存器里的备份。

13. 在GF100芯片上的流多处理器上，寄存器溢出后会将多余的局部数据存放于L1缓存中，由于L1缓存带宽较高，程序仍可保持高性能，这就凸显了L1缓存的重要性。但需要知道，由于寄存器溢出和栈（最多会消耗1KB的L1缓存资源）占用了L1缓存，这将导致其他缓存数据更容易被挤出缓存，因此就会降低缓存的命中率，影响程序性能。

14.C编译器允许main()没有参数，或者有两个参数（有些实现允许更多的参数，但这将是对标准的扩展）。有两个参数时，第一个参数是命令行中的字符串数。按照惯例，这个int参数被称为argc（代表argument count），系统使用空格判断一个字符串结束，另一个字符串开始。第二个参数是一个指向字符串的指针数组。命令行中的每个字符串被存储到内存中，并且分配一个指针指向它。按照惯例，这个指针数组被称为argv(代表argument value)。如果可以（有些操作系统不允许这样做），把程序本身的名字赋值给argv[0],接着，把随后的第一个字符串赋给argv[1]，等等。

  int main(int argc, char **argv) 这种对argv的声明和char *argv[]是等价的。这意味着argv是一个指向“指向字符的指针”的指针。

15.对于结构体，和数组不同的是，一个结构的名字不是该结构的地址，必须使用&运算符。

16.关于编译和链接的一点总结：在编译时，编译器只检测程序语法，和函数，变量是否被声明，如果函数未被声明，编译器会给出一个警告，但是可以生成object file,而在链接程序时，链接器会在所有的Object文件中寻找函数的实现，如果找不到，就会报链接错误码（Linker Error）