cuda学习（linux公社下载地址： http://linux.linuxidc.com/东西比较多）

cuda环境搭建：

1，安装好GPU后，在网上下载CUDA,尽量使用最新的版本吧，这样就可以兼容较新的vs，下载地址：点击打开链接；注意根据自己的系统选择合适的版本。在安转的时候选择自定义安转，尽量安装全面。之后打开vs，在创建中模板中有cuda的模板，这就显示安装成功了。

在C文件中调用内核函数b，要将内核函数b在cu文件中通过其他的函数a调用，之后，c文件中的主函数调用a函数，但是函数a在声明的时候要写成extern"C",在C文件进行调用。

初次编译时会出现的错误：

1，要注意：编写的cuda文件要放在cu文件中，其中的<<<>>>有时会报错，出现的原因就是因为cu文件的编译器不是NVCC，需要在属性中将编译器类型转换一下。

1，cpu中的缓存主要用于减小访存延迟和节约带宽。 缓存在多线程环境下会发生失效反应。在每次线程上下切换止呕，都需要重建缓存上下文。一次缓存失效的代价是几十到上百个时钟周期。同时，为了实现缓存和内存中的数据的一致性，还需要复杂的逻辑进行控制。

2，gpu中则没有复杂的魂村体系和替换极值。GPU的缓存是只读的。因此也不用考虑缓存的一致性问题。GPU缓存的主要功能是用于过滤对存储器控制器的请求，减少对显存的访问。所以缓存的主要功能不是减小访存延迟，而是节约显存带宽，

3，确定了程序中的并行部分，流可以考虑将这部分计算工作交给GPU。运行在ＧＰＵ上的CUDA并行计算函数称为kernel（内核函数)；一个kernel函数并不是一个完整的程序，而是整个cuda程序中的一个可以被并行执行的步骤。一个完整的cuda程序有一系列的设备端kernel函数并行步骤和主机端 的串行处理步骤共同组成。这些语句按照程序中响应的顺序依次执行。

4，在kernel中调用的时候需要尖括号，目的是要将一些参数传递给运行时的系统。这些参数不是传递给设备代码的参数，而是告诉运行时如何启动设备diamante。

5，核函数不能带有返回值，因此返回类型通常是void.

6,cudaMalloc()函数分配设备端的内存；cudaMemcpy()将数据传入或传出设备；cudaFree()释放掉设备中不在使用的内存。

cudaMalloc():1)可以将其分配的指针传递给设备上执行的函数；可以在设备代码中使用其分配的指针进行内存读写操作；可以将其分配的指针传递给在主机上执行的函数；不能再主机代码中使用其分配的指针进行内存读写操作。

主机指针只能访问主机代码中的内存；设置指针也只能访问设备代码中的内存。

使用cudaMemcpy();来访问设备上的内存，

cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice);设备。主机，数据的大小，
cudaMemcpyHostToDevice：从内存复制到显卡内存；cudaMemcpyDeviceToHost - 从显卡内存复制到内存
这个参数告诉运行时源指针是一个设备指针或主机指针，目标指针是一个主机指针或设别指针。

7，size_t::代表内存大小的专用变量类型；cudaError_t：错误处理的专用变量。

8，cuda中出现的HANDLE_ERROR():在book.h中存在。

static void HandleError( cudaError_t err, const char *file, int line ) {
if (err != cudaSuccess) {
printf( "%s in %s at line %d\n", cudaGetErrorString( err ),file, line );
exit( EXIT_FAILURE );
}
}

#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

9，了解系统中有多少个设备支持CUDA 框架，获得cuda设备的数量：使用cudaGetDeviceCount();

10,显示设别类型的结构体：cudaDeviceProp:其中显示设别的相关属性。

11，设备属性的使用：如果核函数与cpu之间需要进行密集的交互，那么就可能需要在集成的ＧＰＵ上运行代码，因为他可以与CPU共享内存。这两个属性可以通过cudaGetDeviceProperties()来查询。

12，memset(&porp，0，sizeof(cudaDeviceProp)):将sizeof()个大小的的0填充到prop中，一般这个函数用来初始化变量。

13，将cudaDeviceProp结构传递给cudaChooseDevice()，这样cuda运行时可以查看是够存在某个设备满足这些条件，cudaChooseDevice()函数将返回一个设备ID，然后我们将这个Id传递给cudaSetDevice(),随后，所有的设备操作豆=都将在这个设备上运行。

14，在函数前面添加关键字：__global__将告诉编译器把该函数放在GPU上运行，通过<<<>>>进行调用。

15，cuda并行编程：

1）在设备上申请空间，避免内存泄漏，在使用完GPU后通过cudaFree()释放他们，通过cudaMemcpy将数据进行在cpu和GPU之间的复制，通过尖括号在主机中的主代码中调用设备上的代码。

2）<<<N,1>>>:这两个参数将传递给运行时，作用是告诉运行时如何启动核函数，第一个参数表示设备在执行核函数时使用的并行线程块的数量。我们将每个并行执行环境称为一个线程块。如果指定kernel<<<256,1>>>():那么将有256个线程块在GPU上运行；了解代码中当前运行的是哪一个线程块：这个问题的答案在于核函数的代码本身。具体来说在于：blockIdx.x:这个变量汇总包含的值就是当前执行设备代码的线程块的索引。使用blockIdx.x：因为cuda支持二维的线程块，使用二维索引可以避免将线性索引转换为矩形索引，使用二维索引比使用一维索引更方便。

3）当启动核函数的时候，我们将并行线程块的数量指定为N,这个并行线程块集合也成为一个线程格，指的是：我们想要一个一维的线程格，其中包含N个线程块，每个线程块中的blockIdx.x值是不同的。第一个为0，最后一个为N-1;  GPU有着完善的内存管理机制，它将结束所有的违反内存访问规则的进程。

4）在启动线程块数组的时候，数组每一维的最大数量都不能超过65535，如果启动的线程块的数量超过了这个限值，那么程序将运行失败。

5)线程中的gridDim是一个常数，用来保存线程格每一维的大小。

6）cuda中的__global__ __ device__    __host__的作用:__global__:表示一个内核函数，有GPU执行；目前的__global__函数必须有cpu调用，并将并行计算任务发送到GPU上执行；__device__:表示这是一个由GPU中的一个线程调用的函数，实际上函数是以__inline形式展开后直接编译到二进制代码中实现的。并不是真正的函数；__host__:和Ｃ或者Ｃ＋＋　中函数一样，是由CPU调用，并执行。

device前缀的函数只能在GPU上运行，所以device修饰的函数不能调用一般的 常见的函数，：global前缀：cuda允许在cpu和GPU上运行，DNA不能运行CPU里常见的函数，Host前缀的函数，是普通函数，默认缺省，可以调用普通函数。

因此，在出现报错如：“ error : calling a __host__ function from a __global__ function is not allowed. ”时候，即为将一个普通的函数错误地添加进入了global前缀定义函数，在CUDA文件.cu文件中是不允许的

7）；在GPU上启动的线程块集合称为一个线程格，线程格既可以是一维的线程块集合，也可以是二维的线程集合，核函数的每个副本都可以通过内置的blockIdx来判断哪个线程正在执行它。可以通过内置的gridDim来获取线程格的大小。

16，线程协作：

1）一些算法需要代码的各个并行副本之间进行通信和协作，在并行执行的cuda c代码段之间实现这种通信很少。

2）核函数的多个并行副本，称为线程块。

3）cuda运行时将这些线程块分解为多个线程，当需要启动多个并行线程块时，只需要将尖括号中的第一个参数由1改为想要启动的线程块数。第二个参数表示cuda运行时在每个线程块中创建的线程数量。总共的启动的线程数量：N个线程块*1个线程/线程块 = N个并行线程。

4）线程块运行和线程运行：a:对线程块的索引是idx = blockIdx.x;一个线程块，对线程索引：int tid = threadIdx.x  ；b；将尖括号中的第二个值进行修改，一个线程块中有几个线程<<<1,N>>>:一个线程块中有N个线程。

5）#define checkCudaErrors(a) do {\
if (cudaSuccess != a) {\
fprintf(stderr, "Cuda runtime error in line%d of file %s:%s\n", __LINE__,__FILE__,cudaGetErrorString(cudaGetLastError));\
exit(EXIT_FAILURE);\
}\
} while (0);

在宏后面加反斜杠的做用：

转载自：http://www.wtoutiao.com/p/K6csca.html

在阅读C语言代码经常可以看到代码中出现反斜杠"\"，不是很明白它的意思，遂对反斜杠"\"的意义和用法进行查询和总结记录。

1. 转义字符

非常常用的一种用法，在反斜杠后面加一个字符，表示一些特定的意思，如：

\n 换行符(LF)

\r 回车(CR) ，相当于键盘上的"Enter"

\t 跳到下一个TAB位置

\0 空字符(NULL)

\' 单引号（撇号）

\" 双引号

\\ 代表一个反斜线字符''\' 等，详细可百度“转义字符”。

例：

a.

printf("Hello World!");

printf("Hello World!");

输出：

Hello World!Hello World!

b.

printf("Hello World!\n");

printf("Hello World!\n");

输出：

Hello World!

Hello World!

2. 继续符

可用于一行的结尾，表示本行与下一行连接起来

C语言中以 ; 作为语句的结束，不以行为单位结束，当一行的内容太长不方便卸载一行时可使用反斜 杠"\"作为继续符

例如：STM32官方库文件"stm32f30x_usart.h"有如下一段：

#define IS_USART_123_PERIPH(PERIPH) (((PERIPH) == USART1) || \

((PERIPH) == USART2) || \

((PERIPH) == USART3))

写成一行意义完全相同：

#define IS_USART_123_PERIPH(PERIPH) (((PERIPH) == USART1) || ((PERIPH) == USART2) || ((PERIPH) == USART3))

（因为显示问题为两行，实际应该为一行）

 

3 宏定义的细节

（1）对于有参数的宏定义，宏定义时，在宏名与带参数的括号之间不应加空格，否则将空格以后的字符都作为替代字符串的一部分。

（2）带参数的宏定义只是进行简单的字符替换，宏展开则是在编译前进行的，在展开时并不分配内存单元，不进行值得传递处理，因此替换不会占用运行时间，只占用编译时间。

（3）宏不存在类型问题，宏名无类型，宏的参数也无类型，只是一个符号代表，展开时代入指定的字符串即可，宏定义时，字符串可以是任何类型的数据。

17，

对于启动核函数时每个线程块中的线程数量，硬件也进行了限制，具体就是：最大的线程数量不能超过设备属性结构中maxThreadsPerBlock域的值。现在许多的图形处理器的限制值是每个线程块是512个线程，对于并行线程长度大于512的矢量进行相加，需要将线程和线程块结合起来才能实现这个计算。

18，这样需要注意的两个地方：核函数的索引计算方法和核函数的调用方法。

19，int  tid = threadIdx.x+blockIdx.x*blockDim.x;  计算多个线程块中总共的线程。

20，内置变量：blockDim,这个变量是一个常数，其中保存的是线程块中每一维的线程数量。cuda运行时允许启动一个二维的线程格。并且线程格中的每个线程块都是一个三维的线程数组。将线程块索引与每个线程块中的线程数量相乘，然后子啊加上线程在线程块中的索引。int offset = x +y*DIM；DIM表示线程块的大小，（即线程的数量），y是线程索引块。x时线程块中的线程索引。tid = threadIdx.x+blockIdx.x*blockDim.x;

21，add<<<(N+127)/128,128>>>（dev_a,dev_b,dev_c）;其中使用N+127保证了开的线程数量>=128,不会开的线程数量不足；在访问输入数组和输出数组之前，必须检查线程的便宜是否位于0到N之间，当索引超过数组的边界时，核函数就会自动的停止执行计算。更重要的是：核函数不会对越过数组边界的内存记性读取或者写入。

22，线程格的每一维的大小不能超过65535,核心是：知道如何计算每个并行线程的初始索引以及如何确定递增的量值。tid = threadIdx.x+blockIdx.x*blockDim.x;线程中每次步长为线程格中正在裕兴的线程数量。这个值等于每个线程块中的线程数量乘以线程格中线程块的数量。为了确保不会启动过多的线程块，我们将线程块的数量固定为某个较小的值。

23，共享内存和同步：

1）将线程块分解为线程的目只是为了解决线程块数量的硬件限制。

2）cuda c支持共享内存，cuda c的关键字__share__添加到变量声明中，这将使这个变量驻留在共享内存中。

3）cuda c编译器将对共享内存中的变量和普通的变量采用不同的处理方式，在GPU上启动的每个线程块，cuda c 编译器将创建该变量的一个副本，线程块中的每个线程都共享这块内存。但线程却无法看到也不能修改其他线程块的变量副本。这就使得一个线程块中的多个线程能够在计算机上进行通信和协作。而且，共享内存的缓冲区驻留在物理GPU上，不是GPU之外的系统内存中，访问共享内存时的延迟要远远低于访问普通缓冲区的延迟。使得共享内存像每个线程块的高速缓存或者中间结果暂存器那样高效。

24，点积运算：（x1,x2,x3,x4）(y1,y2,y3,y4) = x1y1+x2y2+x3y3+x4y4;

1）：首先，可以向矢量加法那样，每个线程将两个相应的元素相乘，然后移动到下两个元素，最终的结果是所有乘积的总和。因此每个线程还要保存它计算的乘积的加和。线程每次对索引的增加值为线程的数量。

2）：volatile;作为指令关键字，确保本条指令不会因编译器的优化而省略。优化器在用到volatile修饰的变量的时候，每次都要重新读取这个变量的值，而不是使用保存在寄存器中的变量。并行设备的硬件寄存器，一个中断服务子程序会访问到非自动变量，多线程应用中被几个任务共享的变量。

一个参数可以是const类型 的，还可以是vilatile:是volatile，因为他可能被意想不到的改变，是const因为程序不应该试图修改它。

3）每个线程块都有共享内存的私有副本。

4）需要确保所有对共享数组cache[]的写入操作在读取cache之前完成，对线程块中的线程进行同步：__syncthreads();这个函数确保线程块中的每个线程都执行完__syncthreads()前面的语句后，才会执行下一条语句。

5）将一个相加过程抽象为更一般的形式，对一个输入数组执行某种计算，然后产生一个更小的结果数组，这种规程称为规约。

6）实现规约：由一个线程在共享内存上进行迭代并计算出总和值，计算的时间和数组的长度成正比。

7）当某些线程需要执行一条指令，而其他的线程不需要执行时，这种情况称为线程发散。在正常的环境中，发散的分支只会使得某些线程处于空闲状态，而其他的线程将执行分支中的代码。

8）使用__syncthreads()情况，线程块中的每个线程都执行了指令，才会执行__syncthreads(),否则不会执行后面的指令。

25，基于共享内存的位图：

1）使用共享内存，和__syncthreads来确保数据在继续进行之前就已经就绪。

2)就是要弄明白线程和位置之间的关系。

26，常量内存和事件（常量内存，和通过事件来测量cuda应用程序的性能）

1）cuda C 程序中可以使用全局内存和共享内存，还支持常量内存。常量内存用于把偶cn在核函数执行期间不会发生变化的数据。在某些情况下，用常量内存来替换全局内存能有效减少内存带宽。

2）常量内存的修饰符：__constant__，

3）在GPU中将数据进行复制到主机：cudaMemcpyToSymbol(s,temp_s,sizeof(*)) :其中s是：__constant__   类 s[size];temp_s是要复制给s的值。最后一个参数是复制数据的大小。

4）使用常量，避免使用cudaMalloc来分配空间。当从主机内存复制到GPU上的常量内存时，我们需要特殊版本的cudaMemcpy();cudaMemcpyToSymbol() 与参数为cudaMemcpyHostToDevice()的cudaMemcpy()之间唯一的差别就是，cudaMemcpyToSymbol()会复制到常量内存中，而cudaMemcpy()会复制到全局内存。

5）常量内存会带来性能的提升，__constant__将把变量的访问限制为只读。和全局内存中读取的数据相比，从常量内存中读取相同的数据可以节约内存带宽。主要由两个原因：a:对常量内存的单次读操作可以广播到其他邻近的线程，这将节约15次读取操作；b;常量内存的数据将缓存起来，因此对相同地址的连续的读操作将不会产生额外的内存通信量。

6）cuda架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“”步调一致”的形式执行，在程序的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

7）当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束。在半线程束中包含了16个线程，基线程束中线程的数量的一半。如果半线程束中的每个线程够从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量的数据，那么这种方式产生的内存流量只是全局的1/16；

8）由于在内存的内容是不会发生变化的，因此硬件将主动把这常量数据缓存在GPU上，在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么将命中缓存，这同样减少了额外的内存流量。

9）半线程束广播：当16个线程读取相同的地址时，这个功能可以极大的提升性能，但当所有16个线程分别读取不同的地址时，他实际会降低性能。读不同的数据，从常量内存读取就慢于从全局内存中读取。

27，使用事件来测量性能：

1）cuda中事件本质上是一个GPU时间，在用户指定的时间点上记录的。获取时间戳：首先创建一个事件，然后记录事件。cudaEvent_t start;

cudaEventCreat(&start);cudaEventRecord(start,0);可以记录cuda运行时记录起始时间和结束时间，

2）防止GPU和cpu在运行时造成逻辑上的错误，需要GPU执行完了在调用cudaEventRecord（）之前的所有语句时，时间才会被记录下来。当GPU完成了之前的工作后，并且记录stop事件后，才能安全的读取stop时间值，告诉CPU在某个事件上同步；使用函数：cudaEventSynchronize(stop)：当GPU执行到stop事件，这才知道stop事件之前的所有的GPU工作已经完成了，因此可以安全的读取stop中保存的时间戳。

记录起始的时间

cudaEvent_t  start stop;

HANDLE_ERROR(cudaEventCreate(&start));

HANDLE_ERROR(cudaEventCreate(&stop));

HANDLE_ERROR(cudaEventRecord(start,0));

//获得结束时间，并显示计时时间
HANDLE_ERROR(cudaEventRecord(stop,0));
HANDLE_ERROR(cudaEventSynchronize(stop));
float elapsedTime;
HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,start,stop));//记录时间差，第一个某浮点变量的地址，之后是两次事件的时间。
printf("time to generate:%3.1fms\n",elapsedTime);


HANDLE_ERROR(cudaEventDestroy(start)); //使用完事件后，要销毁，相当于malloc()分配的内存调用free().
HANDLE_ERROR(cudaEventDestroy(stop));

28,纹理内存：分配和使用纹理内存。一种只读内存。能够提升性能，减少内存流量。

1）纹理内存可以用于通用计算，纹理内存同样缓存在芯片上，在某些情况下，能够减少对内存的请求并提供更高效。

2）纹理缓存是专门为那些在内存访问模式中存在大量的空间局部性的图形应用程序而设计的。纹理缓存为了加速（地址不同，不连续，）这种访问模式而设计的。在这种情况下，使用纹理内存而不使用全局内存，那么性能会提升。

3）需要将输入的数据声明为texture类型的引用，我们使用浮点类型的纹理的引用，因为温度数值是浮点类型的。这些变量将位于GPU上，

texture<float> text;为变量分配了GPU的内存后，需要通过cudaBindTesture()将这些变量绑定到内存缓冲区。这就告诉cuda运行时两件事情：希望将制定的缓冲区作为纹理来使用，希望将纹理引用作为纹理的名字。

4）当读取特殊的函数来告诉GPU将读取的请求转发到纹理内存而不是标准的全局内存，因此，当读取内存时不再使用方括号从缓冲区中读取，而是将blend_kernel()改为tex1Dfetch();

5)全局怒内存和纹理内存差异：tex1Dfetch()是一个编译器的内置函数，由于纹理引用必须声明为文件作用域内的全局变量，因此我们不再将输入缓冲区和输出缓冲区作为参数传递给blend_kernel().因为编译器需要在编译时知道tex1Dfetch()应该对那些纹理采样。不再传递指针，而是将一个布尔标志传递给blend_kernel().这个标志将指明使用哪个缓冲区作为输入，哪个缓冲区作为输出。

6）纹理内存的哟偶是：在芯片上的缓存，在数据访问模式中存在某种空间局部性，

（6,7,8都在将和图形有关的问题，这里主要是数据计算，不是图形的就跳过）

29，原子性（某些情况先回提高性能，但是有时也会降低性能，具体在研究）

1）在编译器中,如果编译器支持的计算功能集版本低于1.1，告诉编译器这个要求时，：nvcc -arch = sm_11: ；在编译需要使用共享内存原子操作的核函数时，使用选项：nvcc -arch = sm_12

2)由于操作的执行过程不能分解为更小的部分，将这种条件限制的操作称为原子操作。当有数千个线程在内存访问上发生竞争时，这些操作能够确保在内存上实现安全的操作。

3）当线程块的数量为GPU中处理器的数量的2倍时，性能最优。

30，流

1）cuda分配主机内存的机制：cudaHostAlloc(),  malloc 和cudaMalloc(),分配的内存之间存在着一个重要的差异，C库函数malloc()将分配标准的，可分配的主机内存，而cudaAlloc()将分配页锁定的主机内存。页锁定内存也称为固定内存，或者不可分页内存。有一个重要属性：操作系统将不会对这块内存分页并交换磁盘，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全的使用某个应用程序访问该内存的物理地址，因为这块内存不会被破坏或者重新定位。

2）GPU知道内存的物理地址，可以直接访问内存技术来在GPU和主机之间复制数据。页锁定内存不能交换到磁盘上，这意味着使用标准的malloc（）调用相比，系统将更快的耗尽内存。

3）cudaHostAlloc((void**)&a,sizeof(*a),cudaHostAllocDefault),最后一个参数的取值范围是一组标识，通过这些标识来修改cudaHostAlloc的行为，并分配不同形式的固定主机内存。释放使用cudaFreeHost().

4)支持设备重叠功能的设备，支持设备重叠功能的GPU能够在执行一个cudaＣ核函数的同时，还能在设备与主机之间执行复制操作，

5）创建流‘

cudaStream_t   stream;  HANDLE_ERROR(cudaStreamCreate(&stream));

6)使用固定内存；1）可以使赋值操作执行的更快，2）将使用一种新的cudaMemcpy()函数，并且在这个新函数中需要页锁定主机内存，

7）将主机和设备上的数据通过异步的方式进行复制：cudaMemcpuAsync(dest,src,size,cudaMemcpyHostToDevice,stream),在流中执行一次内存复制操作，流就像一个有序的工作队列，GPU从该队列中依次去除工作并执行。

8）使用cudaStreamSynchronize()并制定想要等待的流。

31，多GＰＵ系统上的cuda C （暂时用不到，没有细看）

0）固定内存：这种新型的主机内存能够确保不会交换出物理内存，通过cudaHostAlloc()来分配这种内存，并且传递参数cudaHostAllocDefault来获得默认的固定内存。还有其他的标识：cudaHostAllocMapped,,其分配的主机内存也是固定的。可以在cuda C 核函数中直接访问这种类型的主机内存，由于这种内存不需要复制到GPU，因此也称零拷贝内存。

1）cudaHostAllocMapped这个标识告诉运行时将从GPU中访问这块内存，这个标识意味着分配零拷贝内存，对于两个输入缓存区，还指定了标识：cudaHostAllocWriteCombined,这个标识表示：运行时应该将内存分配为合并式写入，内存，这个标识不会改变应用程序的功能，但是可以显著的提升GPU读取内存时的性能，然而，当CPU也要读取这块内存时，合并式写入会显得很低效。因此，在决定是否使用这个标识之前，要考虑应用程序的可能访问模式。

2）cudaHostAlloc()将返回这块内存在CPU上的指针，因此需要调用cudaHostGetDevicePointer()来获得这块内存在GPU上的有效的指针。这些指针将被传递到核函数，

3）在程序中调用cudaThreadSynchronize()将CPU和GPU同步，如果在核函数中会修改零拷贝内存的内容。