CUDA项目配置

（1）打开vs，创建一个空win32程序,即cuda_test项目。

（2）选择cuda_test，点击右键–>生成依赖项–>生成自定义，选择CUDA10.0。

（3）右键源文件文件夹->添加->新建项->选择CUDA C/C++File，取名cuda_main。

（4）点击cuda_main.cu的属性,在配置属性–>常规–>项类型–>选择“CUDA C/C++”。

（5）包含目录配置：

  右键点击项目属性–>属性–>配置属性–>VC++目录–>包含目录

  添加包含目录：$(CUDA_PATH)\include

（6）库目录配置

  VC++目录–>库目录

  添加库目录：$(CUDA_PATH)\lib\x64

（7）依赖项

  配置属性–>链接器–>输入–>附加依赖项

  添加库文件：cublas.lib;cuda.lib;cudadevrt.lib;cudart.lib;cudart_static.lib;OpenCL.lib

 

将CPU及其系统的内存称为主机host,将GPU及其内存称为设备device.

线程块Block由多个线程组成(可以组织为一维、二维和三维），各block是并行执行的，block间无法通信，也没有执行顺序。

线程格Grid由多个线程块组成

线程束Warp：指一个包含32个线程的集合，被“编织在一起”并且“步调一致”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同数据上执行相同的命令。

 

 核函数Kernel:在GPU上执行的函数通常称为核函数，一般通过标识符__global__修饰，调用通过<<<参数1,参数2>>>，用于说明内核函数中的线程数量，以及线程是如何组织的。

CUDA C需要使用某种语法将一个函数标记为“设备代码”，CADA C提供了与C在语言级别上的集成，使得设备调用看起来非常像主机函数调用。尖括号表示要将一些参数传递给运行时系统，告诉运行时如何启动设备代码。

CUDA编程模式

 　　1. 定义需要在 device 端执行的核函数。( 函数声明前加 _golbal_ 关键字 )

       2. 在显存中为待运算的数据以及需要存放结果的变量开辟显存空间。( cudaMalloc 函数实现 )

       3. 将待运算的数据传输进显存。( cudaMemcpy，cublasSetVector 等函数实现 )

       4. 调用 device 端函数，同时要将需要为 device 端函数创建的块数线程数等参数传递进 <<<>>>。( 注: <<<>>>下方编译器可能显示语法错误，不用管 )

       5. 从显存中获取结果变量。( cudaMemcpy，cublasGetVector 等函数实现 )

       6. 释放申请的显存空间。( cudaFree 实现 )

函数声明

　　　　1. __device__

　　　　表明此函数只能在 GPU 中被调用，在 GPU 中执行。这类函数只能被 __global__ 类型函数或 __device__ 类型函数调用。

　　　　2. __global__

　　　　表明此函数在 CPU 上调用，在 GPU 中执行。这也是以后会常提到的 "内核函数"，有时为了便于理解也称 "device" 端函数。

　　　　3. __host__

　　　　表明此函数在 CPU 上调用和执行，这也是默认情况。

　　内核函数配置运算符 <<<>>> - 这个运算符在调用内核函数的时候使用，一般情况下传递进三个参数：

　　　　1. 块数

　　　　2. 线程数

　　　　3. 共享内存大小 (此参数默认为0 )

几个内置变量

1. 　　 threadIdx，顾名思义获取线程thread的ID索引；如果线程是一维的那么就取threadIdx.x，二维的还可以多取到一个值threadIdx.y，以此类推到三维 threadIdx.z。
2.         blockIdx，线程块的ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。
3.       blockDim，线程块的维度，同样有blockDim.x，blockDim.y，blockDim.z。
4.       gridDim，线程格的维度，同样有gridDim.x，gridDim.y，gridDim.z。

        5. 对于一维的block，线程的threadID=threadIdx.x。

        6. 对于大小为（blockDim.x, blockDim.y）的 二维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x。

        7. 对于大小为（blockDim.x, blockDim.y, blockDim.z）的 三维 block，线程的threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。

        8. 对于计算线程索引偏移增量为已启动线程的总数。如stride = blockDim.x * gridDim.x; threadId += stride。

GPU内存

全局内存

通俗意义上的设备内存

共享内存

使用__shared__关键字声明，例如__shared__ float cache[10]，对于GPU上启动的每个线程块，CUDA C编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程能够在计算上通信和协作。

常量内存

使用关键字__constant__声明，为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。常量内存用于保存在核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。当我们需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。

从常量内存中读取数据可以节约内存带宽，主要有两个原因：

1. 对常量内存的单次读操作可以广播到邻近线程，这将节约约15次读取操作
2. 常量内存的数据将缓存起来，因此对相同地址的连续操作将不会产生额外的内存通信量

纹理内存

 固定内存

 

常用函数

cudaGetDeviceCount()　　获取显示设备数目

cudaGetDeviceProperties()　　获取设备属性

cudaChooseDevice()　　根据指定的属性条件选择设备

cudaSetDevice()　　指定使用的显示设备

cudaMalloc()　　在设备中分配空间

cudaMemcpy()　　host和device之间拷贝内存，

cudaFree()　　释放显存

__syncthreads()　　用于同一线程块内线程间的同步，__syncthreads() is you garden variety thread barrier. Any thread reaching the barrier waits until all of the other threads in that block also reach it.

 矢量求和

1 #include "cuda_runtime.h" 2 #include "cuda.h" 3 #include "device_launch_parameters.h" 4  5 #include 
      
        6 #include 
       
         7 #define N 10 8  9 __global__ void add(int* a, int* b, int*c)10 {11     int tid = blockIdx.x;12     if (tid < N)13     {14         c[tid] = a[tid] + b[tid];15     }16 }17 18 int main(void)19 {20     int a[N], b[N], c[N];21     int *dev_a, *dev_b, *dev_c;22     cudaMalloc((void**)&dev_a, N * sizeof(int));23     cudaMalloc((void**)&dev_b, N * sizeof(int));24     cudaMalloc((void**)&dev_c, N * sizeof(int));25     for (int i = 0; i < N; ++i)26     {27         a[i] = -i;28         b[i] = i*i;29     }30     cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);31     cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);32     add << 
        
         > > (dev_a, dev_b, dev_c);33     cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);34     for (int i = 0; i < N; ++i)35     {36         printf("%d+%d=%d\n", a[i], b[i], c[i]);37     }38     cudaFree(dev_a);39     cudaFree(dev_b);40     cudaFree(dev_c);41 42     getchar();43     return 0;44 }
        
       
      

 调用核函数<<<>>>中

第一个参数表示设备在执行核函数时使用的并行线程块数量，即创建多少个核函数的副本并以并行的方式执行它们。内置变量blockIdx包含的值就是当前执行设备代码的线程块的索引。硬件限制线程块数量不能超过65535，内置变量blockDim保存的是三维的线程块中线程的维度。即CUDA运行时允许启动一个二维线程格，且线程格中的每个线程块都是一个三维的线程数组。

第二个参数表示CUDA运行时在每个线程块中创建的线程数量，内置参数threadIdx为线程索引。硬件限制每个线程块中线程数量不能超过设备属性结构中maxThreadsPerBlock的值。

并行线程块集合也称为一个线程格Grid。

事件

cuda中的事件本质上是一个GPU时间戳

1 /*  2 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.  3 *  4 * NVIDIA Corporation and its licensors retain all intellectual property and  5 * proprietary rights in and to this software and related documentation.  6 * Any use, reproduction, disclosure, or distribution of this software  7 * and related documentation without an express license agreement from  8 * NVIDIA Corporation is strictly prohibited.  9 * 10 * Please refer to the applicable NVIDIA end user license agreement (EULA) 11 * associated with this source code for terms and conditions that govern 12 * your use of this NVIDIA software. 13 * 14 */ 15  16  17 #include "cuda.h" 18 #include "../common/book.h" 19 #include "../common/cpu_bitmap.h" 20  21 #define DIM 1024 22  23 #define rnd( x ) (x * rand() / RAND_MAX) 24 #define INF     2e10f 25  26 struct Sphere { 27     float   r, b, g; 28     float   radius; 29     float   x, y, z; 30     __device__ float hit(float ox, float oy, float *n) { 31         float dx = ox - x; 32         float dy = oy - y; 33         if (dx*dx + dy*dy < radius*radius) { 34             float dz = sqrtf(radius*radius - dx*dx - dy*dy); 35             *n = dz / sqrtf(radius * radius); 36             return dz + z; 37         } 38         return -INF; 39     } 40 }; 41 #define SPHERES 200 42  43 __constant__ Sphere s[SPHERES]; 44  45 __global__ void kernel(unsigned char *ptr) { 46     // map from threadIdx/BlockIdx to pixel position 47     int x = threadIdx.x + blockIdx.x * blockDim.x; 48     int y = threadIdx.y + blockIdx.y * blockDim.y; 49     int offset = x + y * blockDim.x * gridDim.x; 50     float   ox = (x - DIM / 2); 51     float   oy = (y - DIM / 2); 52  53     float   r = 0, g = 0, b = 0; 54     float   maxz = -INF; 55     for (int i = 0; i
      
        maxz) { 59             float fscale = n; 60             r = s[i].r * fscale; 61             g = s[i].g * fscale; 62             b = s[i].b * fscale; 63             maxz = t; 64         } 65     } 66  67     ptr[offset * 4 + 0] = (int)(r * 255); 68     ptr[offset * 4 + 1] = (int)(g * 255); 69     ptr[offset * 4 + 2] = (int)(b * 255); 70     ptr[offset * 4 + 3] = 255; 71 } 72  73 // globals needed by the update routine 74 struct DataBlock { 75     unsigned char   *dev_bitmap; 76 }; 77  78 int main(void) { 79     DataBlock   data; 80     // capture the start time 81     cudaEvent_t     start, stop; 82     HANDLE_ERROR(cudaEventCreate(&start)); 83     HANDLE_ERROR(cudaEventCreate(&stop)); 84     HANDLE_ERROR(cudaEventRecord(start, 0)); 85  86     CPUBitmap bitmap(DIM, DIM, &data); 87     unsigned char   *dev_bitmap; 88  89     // allocate memory on the GPU for the output bitmap 90     HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, 91         bitmap.image_size())); 92  93     // allocate temp memory, initialize it, copy to constant 94     // memory on the GPU, then free our temp memory 95     Sphere *temp_s = (Sphere*)malloc(sizeof(Sphere) * SPHERES); 96     for (int i = 0; i
       
        > >(dev_bitmap);113 114     // copy our bitmap back from the GPU for display115     HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap,116         bitmap.image_size(),117         cudaMemcpyDeviceToHost));118 119     // get stop time, and display the timing results120     HANDLE_ERROR(cudaEventRecord(stop, 0));121     HANDLE_ERROR(cudaEventSynchronize(stop));122     float   elapsedTime;123     HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,124         start, stop));125     printf("Time to generate:  %3.1f ms\n", elapsedTime);126 127     HANDLE_ERROR(cudaEventDestroy(start));128     HANDLE_ERROR(cudaEventDestroy(stop));129 130     HANDLE_ERROR(cudaFree(dev_bitmap));131 132     // display133     bitmap.display_and_exit();134 }
       
      

由于cuda事件是直接在GPU上实现的，因此它们不适用于对同时包含设备代码和主机代码的混合代码计时，也就是说，如果试图通过cuda事件对核函数和设备内存复制之外的代码进行计时，将会得到不可靠的结果。