cuda基础
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 )
几个内置变量
- threadIdx，顾名思义获取线程thread的ID索引；如果线程是一维的那么就取threadIdx.x，二维的还可以多取到一个值threadIdx.y，以此类推到三维 threadIdx.z。
- blockIdx，线程块的ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。
- blockDim，线程块的维度，同样有blockDim.x，blockDim.y，blockDim.z。
- 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()会复制到全局内存。
从常量内存中读取数据可以节约内存带宽，主要有两个原因：
- 对常量内存的单次读操作可以广播到邻近线程，这将节约约15次读取操作
- 常量内存的数据将缓存起来，因此对相同地址的连续操作将不会产生额外的内存通信量
纹理内存
固定内存
常用函数
cudaGetDeviceCount() 获取显示设备数目
cudaGetDeviceProperties() 获取设备属性
cudaChooseDevice() 根据指定的属性条件选择设备
cudaSetDevice() 指定使用的显示设备
cudaMalloc() 在设备中分配空间
cudaMemcpy() host和device之间以同步方式拷贝内存
cudaMemcpyAsync() host和device之间以异步方式拷贝内存，任何传递给cudaMemcpyAsync()的主机内存指针都必须已通过cudaHostAlloc()分配好内存。
cudaMemset() Initializes or sets device memory to a value
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.
atomicAdd() 原子加
cudaThreadSynchronize() Wait for compute device to finish
矢量求和
1 #include "cuda_runtime.h" 2 #include "cuda.h" 3 #include "device_launch_parameters.h" 4 5 #include <iostream> 6 #include <cstdlib> 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 << <N, 1 >> > (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<SPHERES; i++) { 56 float n; 57 float t = s[i].hit(ox, oy, &n); 58 if (t > 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<SPHERES; i++) { 97 temp_s[i].r = rnd(1.0f); 98 temp_s[i].g = rnd(1.0f); 99 temp_s[i].b = rnd(1.0f); 100 temp_s[i].x = rnd(1000.0f) - 500; 101 temp_s[i].y = rnd(1000.0f) - 500; 102 temp_s[i].z = rnd(1000.0f) - 500; 103 temp_s[i].radius = rnd(100.0f) + 20; 104 } 105 HANDLE_ERROR(cudaMemcpyToSymbol(s, temp_s, 106 sizeof(Sphere) * SPHERES)); 107 free(temp_s); 108 109 // generate a bitmap from our sphere data 110 dim3 grids(DIM / 16, DIM / 16); 111 dim3 threads(16, 16); 112 kernel << <grids, threads >> >(dev_bitmap); 113 114 // copy our bitmap back from the GPU for display 115 HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap, 116 bitmap.image_size(), 117 cudaMemcpyDeviceToHost)); 118 119 // get stop time, and display the timing results 120 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 // display 133 bitmap.display_and_exit(); 134 }
由于cuda事件是直接在GPU上实现的，因此它们不适用于对同时包含设备代码和主机代码的混合代码计时，也就是说，如果试图通过cuda事件对核函数和设备内存复制之外的代码进行计时，将会得到不可靠的结果。
页锁定主机内存
通常通过cudaMalloc()在GPU上分配内存，通过标准的C库函数malloc()在主机上分配内存，另外cuda运行时还提供了自己独有的机制来分配主机内存：cudaHostAlloc()。C库函数将分配标准的，可分页的主机内存，而cudaHostAlloc()将分配叶锁定的主机内存，操作系统不会对这块内存分页交换到磁盘上，从而确保了该内存始终驻留在物理内存中，因此操作系统能够安全的使某个应用程序访问该内存的物理地址，因为这块内存不会被破坏或者重定位。由于GPU知道内存的物理地址，因此可以通过“直接内存访问DMA”技术来在GPU和主机之间访问复制数据，无需CPU的介入。通过cudaFreeHost()释放由cudaHostAlloc()分配的内存。
流
cuda流表示一个GPU操作队列，并且该操作队列中的操作将以指定的顺序执行。可以将每个流视为GPU上的一个任务，并且这些任务可以并行执行。
cudaStream_t stream; cudaStreamCreate(&stream); //创建流
cudaStreamSynchronize(stream)； //等待流任务完成
cudaStreamDestroy(stream)；//Destroys and cleans up an asynchronous stream
在任何支持内存复制和核函数执行相互重叠的设备上，当使用多个流时，应用程序的性能都会得到提升。
GPU硬件中并没有流的概念，而是包含一个或多个引擎来执行内存复制操作，以及一个引擎来执行核函数。CUDA驱动程序负责按照操作顺序把它们调度到硬件上执行，从而维持流内部的依赖性。
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 "../common/book.h" 18 19 #define N (1024*1024) 20 #define FULL_DATA_SIZE (N*20) 21 22 23 __global__ void kernel( int *a, int *b, int *c ) { 24 int idx = threadIdx.x + blockIdx.x * blockDim.x; 25 if (idx < N) { 26 int idx1 = (idx + 1) % 256; 27 int idx2 = (idx + 2) % 256; 28 float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f; 29 float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f; 30 c[idx] = (as + bs) / 2; 31 } 32 } 33 34 35 int main( void ) { 36 cudaDeviceProp prop; 37 int whichDevice; 38 HANDLE_ERROR( cudaGetDevice( &whichDevice ) ); 39 HANDLE_ERROR( cudaGetDeviceProperties( &prop, whichDevice ) ); 40 if (!prop.deviceOverlap) { 41 printf( "Device will not handle overlaps, so no speed up from streams\n" ); 42 return 0; 43 } 44 45 cudaEvent_t start, stop; 46 float elapsedTime; 47 48 cudaStream_t stream0, stream1; 49 int *host_a, *host_b, *host_c; 50 int *dev_a0, *dev_b0, *dev_c0; 51 int *dev_a1, *dev_b1, *dev_c1; 52 53 // start the timers 54 HANDLE_ERROR( cudaEventCreate( &start ) ); 55 HANDLE_ERROR( cudaEventCreate( &stop ) ); 56 57 // initialize the streams 58 HANDLE_ERROR( cudaStreamCreate( &stream0 ) ); 59 HANDLE_ERROR( cudaStreamCreate( &stream1 ) ); 60 61 // allocate the memory on the GPU 62 HANDLE_ERROR( cudaMalloc( (void**)&dev_a0, 63 N * sizeof(int) ) ); 64 HANDLE_ERROR( cudaMalloc( (void**)&dev_b0, 65 N * sizeof(int) ) ); 66 HANDLE_ERROR( cudaMalloc( (void**)&dev_c0, 67 N * sizeof(int) ) ); 68 HANDLE_ERROR( cudaMalloc( (void**)&dev_a1, 69 N * sizeof(int) ) ); 70 HANDLE_ERROR( cudaMalloc( (void**)&dev_b1, 71 N * sizeof(int) ) ); 72 HANDLE_ERROR( cudaMalloc( (void**)&dev_c1, 73 N * sizeof(int) ) ); 74 75 // allocate host locked memory, used to stream 76 HANDLE_ERROR( cudaHostAlloc( (void**)&host_a, 77 FULL_DATA_SIZE * sizeof(int), 78 cudaHostAllocDefault ) ); 79 HANDLE_ERROR( cudaHostAlloc( (void**)&host_b, 80 FULL_DATA_SIZE * sizeof(int), 81 cudaHostAllocDefault ) ); 82 HANDLE_ERROR( cudaHostAlloc( (void**)&host_c, 83 FULL_DATA_SIZE * sizeof(int), 84 cudaHostAllocDefault ) ); 85 86 for (int i=0; i<FULL_DATA_SIZE; i++) { 87 host_a[i] = rand(); 88 host_b[i] = rand(); 89 } 90 91 HANDLE_ERROR( cudaEventRecord( start, 0 ) ); 92 // now loop over full data, in bite-sized chunks 93 for (int i=0; i<FULL_DATA_SIZE; i+= N*2) { 94 // enqueue copies of a in stream0 and stream1 95 HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i, 96 N * sizeof(int), 97 cudaMemcpyHostToDevice, 98 stream0 ) ); 99 HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N, 100 N * sizeof(int), 101 cudaMemcpyHostToDevice, 102 stream1 ) ); 103 // enqueue copies of b in stream0 and stream1 104 HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i, 105 N * sizeof(int), 106 cudaMemcpyHostToDevice, 107 stream0 ) ); 108 HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N, 109 N * sizeof(int), 110 cudaMemcpyHostToDevice, 111 stream1 ) ); 112 113 // enqueue kernels in stream0 and stream1 114 kernel<<<N/256,256,0,stream0>>>( dev_a0, dev_b0, dev_c0 ); 115 kernel<<<N/256,256,0,stream1>>>( dev_a1, dev_b1, dev_c1 ); 116 117 // enqueue copies of c from device to locked memory 118 HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0, 119 N * sizeof(int), 120 cudaMemcpyDeviceToHost, 121 stream0 ) ); 122 HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1, 123 N * sizeof(int), 124 cudaMemcpyDeviceToHost, 125 stream1 ) ); 126 } 127 HANDLE_ERROR( cudaStreamSynchronize( stream0 ) ); 128 HANDLE_ERROR( cudaStreamSynchronize( stream1 ) ); 129 130 HANDLE_ERROR( cudaEventRecord( stop, 0 ) ); 131 132 HANDLE_ERROR( cudaEventSynchronize( stop ) ); 133 HANDLE_ERROR( cudaEventElapsedTime( &elapsedTime, 134 start, stop ) ); 135 printf( "Time taken: %3.1f ms\n", elapsedTime ); 136 137 // cleanup the streams and memory 138 HANDLE_ERROR( cudaFreeHost( host_a ) ); 139 HANDLE_ERROR( cudaFreeHost( host_b ) ); 140 HANDLE_ERROR( cudaFreeHost( host_c ) ); 141 HANDLE_ERROR( cudaFree( dev_a0 ) ); 142 HANDLE_ERROR( cudaFree( dev_b0 ) ); 143 HANDLE_ERROR( cudaFree( dev_c0 ) ); 144 HANDLE_ERROR( cudaFree( dev_a1 ) ); 145 HANDLE_ERROR( cudaFree( dev_b1 ) ); 146 HANDLE_ERROR( cudaFree( dev_c1 ) ); 147 HANDLE_ERROR( cudaStreamDestroy( stream0 ) ); 148 HANDLE_ERROR( cudaStreamDestroy( stream1 ) ); 149 150 return 0; 151 }
零拷贝主机内存
通过cudaHostAlloc()分配叶锁定内存时，如果参数标志指定为cudaHostAllocMapped时，这种内存除了可以用于主机和设备之间的内存复制外，还可以在CUDA C 核函数中直接访问这种类型的主机内存，由于这种内存不需要复制到GPU，所以称为零复制内存。
通过cudaHostGetDevicePointer()来获取这块主机内存在GPU上的有效指针。
通过查看属性canMapHostMemory值可查看硬件是否支持映射主机内存。
通过cudaSetDeviceFlags(cudaDeviceMapHost)；将运行时设置为能分配零拷贝内存的状态。
There is a limit to the number of threads per block,On current GPUs,a thread block may contain up to 1024 threads.
Each thread has private local memory. Each thread block has shared memory visible to all threads of the block and with the same lifetime as the block. All threads have access to the same global memory.
The compute capability of a device is represented by a version number, also sometimes called its "SM version". This version number identifies the features supported by the GPU hardware and is used by applications at runtime to determine which hardware features and/or instructions are available on the present GPU. Devices with the same major revision number are of the same core architecture. The major revision number is 7 for devices based on the Volta architecture, 6 for devices based on the Pascal architecture, 5 for devices based on the Maxwell architecture, 3 for devices based on the Kepler architecture, 2 for devices based on the Fermi architecture, and 1 for devices based on the Tesla architecture.
The cuda runtime is implemented in the cudart library,which is linked to the application.
