第一章
- 用途：加速三维成像，计算流体力学，分子模拟计算等等，这些应用场景大都需要进行大量的科学计算，GPU的强大并行计算的能力使得其能够完成这些超大量的计算任务，为任务缩短计算时间，降低计算代价；
第二章
- 首先需要安装cuda toolkit，添加环境变量；
- 其次，vs创建工程，并在项目中加入cuda的bin，inlcude和lib；
- 编写.cu代码，需要以__global__开头，并且使用extern "C"声明，在cpp文件中同样需要使用extern "C"声明，然后再调用；
第三章
3.1 and 3.2
-
第一个cuda程序：首先，创建.cu文件，添加如下代码：
#include "cuda_runtime.h" #include "device_launch_parameters.h" __global__ void kernel() {} extern "C" void test_kernel(){ kernel <<<1, 1 >>> (); }然后创建.cpp文件，添加如下代码：
#include <iostream> // test1 extern "C" void test_kernel(); int main() { //test1 test_kernel(); printf("hello, world!\n"); return 0; }cuda代码都是以__global__开头，这是告诉编译器，该函数应该编译为在设备上运行而非主机上，即该符号将函数kernel()标记为设备代码；在该实例中，函数kernel()将被交给编译设备代码的编译器，而main()函数则是交给主机编译器；而其中，调用时使用的尖括号<<<>>>表示将一些参数传递给运行时系统，这些参数并不是传递给设备代码的参数，这是告诉运行时系统如何启动设备代码，传递给设备代码本省的参数是放在圆括号中传递的，如同C语言中普通函数的参数传递一样。
需要注意的是，在.cpp中使用__global__编写cuda代码时，首先需要#include “cuda_runtime.h” 和#include “device_launch_parameters.h”；其次，需要将.cpp文件的属性中项目类型选为CUDA C/C++，否则将无法识别尖括号；
注意：程序员一定不能在在主机代码中对cudaMalloc()返回的指针进行解引用，主机代码可以将这个指针作为参数出传递，对其执行算术运算，还可以将其转换为另一种类型，但绝不能使用这个指针来读取或者写入内存；
-
设备指针的使用限制：
- 可以将cudaMalloc()分配的指针传递给设备上执行的函数；
- 可以在设备代码中使用cudaMalloc()分配的指针进行内存的读写；
- 可以将cudaMalloc()分配的指针传递给主机上执行的函数；但最终都是传递给设备代码的；
- 不能在主机代码中使用cudaMalloc()分配的指针进行读写；
- 不能使用标准的 C 的free()来释放cudaMalloc()分配的内存，需要使用cudaFree()；
- 如果需要访问设备上的内存，需要使用cudaMemcpy来拷贝到主机内存进行访问，其最后一个参数为cudaMemcpyDeviceToHost，其告诉运行时系统源指针是一个设备指针，目标指针是一个主机指针；cudaMemcpyHostToDevice则相反， cudaMemcpyDeviceToDevice表示两个指针都在设备上；
-
polo
3.3 查询设备
-
可以通过cudaGetDeviceCount()查询cuda设备的数量，同时可以根据其设备编号使用cudaGetDeviceProp()查询设备的属性，代码如下：
#include <iostream> #include<cuda_runtime.h> __global__ void kernel(){} int main() { int count; cudaGetDeviceCount(&count); std::cout << " count = " << count << std::endl; cudaDeviceProp prop; for (int i = 0; i < count; ++i) { cudaGetDeviceProperties(&prop, i); std::cout << "prop.name: " << prop.name << std::endl; std::cout << "compute capability: " << prop.major <<" "<< prop.minor << std::endl; } //std::cout << "Hello World!\n"; } -
但是对每个设备进行迭代查询满足要求的设备比较繁琐，可以使用cuda运行时的一种自动的方式来执行这个迭代操作
#include <iostream> #include<cuda_runtime.h> __global__ void kernel(){} int main() { cudaDeviceProp prop; int dev; cudaGetDevice(&dev); std::cout << "id of current cuda device: " << dev << std::endl; //填充cudaDeviceProp结构 memset(&prop, 0, sizeof(cudaDeviceProp)); prop.major = 7; prop.minor = 5; //传递给cudaChooseDevice cudaChooseDevice(&dev, &prop); std::cout << "id of current cuda device closest to 7.5: " << dev << std::endl; cudaSetDevice(dev); //std::cout << "Hello World!\n"; } -
polo
第四章
-
矢量求和运算
-
cpu上的实现
#include <iostream> #include<cuda_runtime.h> #define N 10 void add_cpu(int *a, int *b, int *c) { int tid = 0; while (tid < N) { c[tid] = a[tid] + b[tid]; tid += 1; } } int main() { int a[N], b[N], c[N]; //在cpu上为数组a和b赋值 for (int i = 0; i < N; ++i) { a[i] = -i; b[i] = i * i; } add_cpu(a, b, c); //结果 for (int i = 0; i < N; ++i) { std::cout << "a[" << i << "] +b[" << i << "] = " << c[i] << std::endl; } return 0; }如果使用cpu来完成加速计算，则需要充分利用多线程来进行计算，使用不同的内核来计算不同位置的球和操作，这样使得程序代码变得复杂。
-
gpu上的实现
#include <iostream> #include<cuda_runtime.h> #include<device_launch_parameters.h> #define N 10 void add_cpu(int *a, int *b, int *c) { int tid = 0; while (tid < N) { c[tid] = a[tid] + b[tid]; tid += 1; } } __global__ void add(int *a, int *b, int *c) { int tid = blockIdx.x; if (tid < N) { c[tid] = a[tid] + b[tid]; } } int main() { int a[N], b[N], c[N]; int *dev_a, *dev_b, *dev_c; //GPU上分配内存 cudaMalloc((void**)&dev_a, N * sizeof(int)); cudaMalloc((void**)&dev_b, N * sizeof(int)); cudaMalloc((void**)&dev_c, N * sizeof(int)); //在cpu上为数组a和b赋值 for (int i = 0; i < N; ++i) { a[i] = -i; b[i] = i * i; } //将数据拷贝到gpu cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice); add << <N, 1 >> > (dev_a, dev_b, dev_c); //dev_c从Gpu复制到cpu cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost); //结果 for (int i = 0; i < N; ++i) { std::cout << "a[" << i << "] +b[" << i << "] = " << c[i] << std::endl; } //释放cuda显存 cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; }-
gpu上实现的过程大致和cpu上相同。需要注意的就是gpu上创建内存不能再cpu中写，cpu上的不能再gpu中写。因此多了很多的cudaMemcpy的操作，同时也使用cudaFree来释放显存；
-
尖括号中的参数含义：add<<<N, 1>>>(param1, param2,…)
第一个参数N表示设备在执行核函数时使用的并行线程块的数量，可以理解为运行时系统创建了N个和函数的副本，并以并行方式运行他们。每个并行执行环境都称为一个线程块，此例中有N个线程块；
-
如何知道运行的是哪一个线程块：答案是通过blocIdx.x实现。blockIdx为一个内置变量，在cuda运行时已经定义了该变量，位于device_launch_parameters.h中；由于cuda支持二维的线程块数组，对于二维空间的计算问题，如图像的处理；
-
每个线程块的blockIdx.x的值是不相同的，需要注意的是tid<N，否则容易发生非法访问内存；
-
需要注意的是，线程块数组中每一维的最大数量都不能超过65535；
-
polo
-
-
polo
-
-
polo
第五章 线程协作
-
并行线程块的分解：cuda运行时系统将线程块分解为多个线程。add<<<N,1>>>则表示N个线程块（N个线程块可以并行），每个线程块中拥有一个线程。
-
同样的矢量相加的任务我们使用并行的线程去执行，即add<<<1, N>>>。虽然这样做相比于使用多个线程块没有任何优势，但其能完成线程块无法完成的工作。代码如下：
// CUDATest.cpp : 此文件包含 "main" 函数。程序执行将在此处开始并结束。 // #include <iostream> #include<cuda_runtime.h> #include<device_launch_parameters.h> #define N 10 void add_cpu(int *a, int *b, int *c) { int tid = 0; while (tid < N) { c[tid] = a[tid] + b[tid]; tid += 1; } } __global__ void add(int *a, int *b, int *c) { int tid = threadIdx.x; if (tid < N) { c[tid] = a[tid] + b[tid]; } } int main() { int a[N], b[N], c[N]; int *dev_a, *dev_b, *dev_c; //GPU上分配内存 cudaMalloc((void**)&dev_a, N * sizeof(int)); cudaMalloc((void**)&dev_b, N * sizeof(int)); cudaMalloc((void**)&dev_c, N * sizeof(int)); //在cpu上为数组a和b赋值 for (int i = 0; i < N; ++i) { a[i] = -i; b[i] = i * i; } //将数据拷贝到gpu cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice); add << <1, N >> > (dev_a, dev_b, dev_c); //dev_c从Gpu复制到cpu cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost); //结果 for (int i = 0; i < N; ++i) { std::cout << "a[" << i << "] +b[" << i << "] = " << c[i] << std::endl; } //释放cuda显存 cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; }可以看出，只有两处不同，即threadIdx.x和add<<<1, N>>>。
-
硬件设备限制了线程块数量不能超过65535，而每个线程块中线程的数量限制不能超过512。如果使用1个线程块去对超过512的矢量进行相加，则需要结合线程块和线程才能实现。需要改动的地方有两个，即核函数中的索引计算方法和核函数的调用方式。
-
多个线程块包含多个线程
-
索引方式的改变：int tid = threadIdx.x +blockIdx.x*blockDim.x。对于所有的线程块来说，blockDim是一个常数，表示线程块中每一维的线程数量；
-
核函数的调用：当对任意长度的矢量进行求和时，可能矢量的长度超过了总的线程数=线程块数*线程块内线程数，这导致需要寻找其他的方案实现。方案就是将线程块也进行类似cpu多核处理的操作，即将tid进行自增，代码如下：
__global__ void add(int *a, int *b, int *c){ int tid = threadIdx.x + blockIdx.x * blockDim.x; while(tid < N){ c[tid] = a[tid] + b[tid]; tid += blockIdx.x * gridDim.x; //递增的是当前线程格正在运行的线程数量 } }上数修改意味着每次递增的数为线程块数*格子数。因为每个grid容纳的线程块数为gridDim.x，每个线程块的线程个数为blockIdx.x。我们只需要知道每个并行线程的初始化索引，以及如何确定递增的量值即可，所以需要线程索引和线程块索引进行线性化，即首先找到线程块索引，然后找到线程索引，代码为：
int tid = threadIdx.x + blockIdx.x * blockDim.x;每个线程计算完当前索引上的任务之后，需要对索引进行递增，其递增的步长为线程格中正在运行的线程数量，其数量为每个线程块中的线程数量乘以线程格中线程块的数量。于是递增代码为:
tid += blockIdx.x * gridDim.x; //递增的是当前线程格正在运行的线程数量,blockIdx.x为每个线程块中的线程数，gridDim.x为线程格中的线程块数 -
polo
-
-
在GPU上使用线程实现波纹的效果
- 略
-
共享内存和同步
-
线程块分解为线程完全可以由cuda运行时在幕后完成。但还有其他的重要原因需要将线程块分解为多个线程。使用__share__关键字可以将申明的变量驻留在内存中；
-
共享内存缓冲区驻留在物理GPU上，而非GPU之外的内存。在GPU上每启动一个线程块，cuda c编译器都将创建一个该变量的一个副本，线程块中的每个变量都共享这块内存，这就使得一个线程块中的多个线程可以进行通讯、协作；
-
有共享，就有竞争。和cpu中的多线程一样，gou中共享内存也需要设置竞争策略，如何及进行同步是一个问题？
-
核函数调用方式：线程块的大小设为固定值，即单个线程块中线程的数量固定，如128，则可以启动N/128个线程块。一般N/128需要向上取整，避免N小于线程块固定大小值而无法启动线程；实现为(N+127)/128，代码为：add<<<(N+127)/128, 128>>>(dev_a, dev_b, dev_c)；
共享内存的代码为：
// CUDATest.cpp : 此文件包含 "main" 函数。程序执行将在此处开始并结束。 // #include <iostream> #include<cuda_runtime.h> #include<device_launch_parameters.h> #define imin(a, b)(a<b?a:b) const int N = 33 * 1024; const int threadsPerBlock = 256; const int blocksPerGrid = imin(32, (N + threadsPerBlock - 1) / threadsPerBlock); __global__ void dot(float *a, float *b, float *c) { __shared__ float cache[threadsPerBlock];//内存缓冲区 int tid = threadIdx.x + blockIdx.x * blockDim.x; //blockIdx.x 表示block中的线程数，blockDim.x表示线程块索引 int cacheIndex = threadIdx.x; float temp = 0; while (tid < N) { temp = a[tid] * b[tid]; tid += blockIdx.x * gridDim.x;//线程格中的线程块数 * 线程块中的线程数 = 当前线程格正在运行的线程数 } cache[cacheIndex] = temp; __syncthreads(); //线程块中的线程同步 //归约 int i = blockDim.x / 2; while (i != 0) { if (cacheIndex < i) { cache[cacheIndex] += cache[cacheIndex + i]; } __synctheads(); i /= 2; } if (cacheIndex == 0) {//可以选择任意一个线程将cache[0]写入全局内存 c[blockIdx.x] = cache[0]; } } int main() { float *a, *b, c, *partial_c; float *dev_a, *dev_b, *dev_partial_c; //cpu上分配内存 //a = (float*)malloc(N * sizeof(float)); //b = (float*)malloc(N * sizeof(float)); //partial_c = (float*)malloc(blocksPerGrid * sizeof(float)); a = new float[N]; b = new float[N]; partial_c = new float[blocksPerGrid]; cudaMalloc((void**)&dev_a, N * sizeof(float)); cudaMalloc((void**)&dev_b, N * sizeof(float)); cudaMalloc((void**)&dev_partial_c, blocksPerGrid * sizeof(float)); //填充主机内存 for (int i = 0; i < N; ++i) { a[i] = i; b[i] = i * 2; } //将数据拷贝到gpu cudaMemcpy(dev_a, a, N * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * sizeof(float), cudaMemcpyHostToDevice); dot << <blocksPerGrid, threadsPerBlock >> > (dev_a, dev_b, dev_partial_c); //dev_c从Gpu复制到cpu cudaMemcpy(partial_c, dev_partial_c, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost); //结果 for (int i = 0; i < N; ++i) { std::cout << "a[" << i << "] +b[" << i << "] = " << partial_c[i] << std::endl; } //释放cuda显存 cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_partial_c); //释放cpu上的内存 delete[] a; delete [] b; delete[] partial_c; return 0; } -
共享内存指的是同一个线程块中的线程共享一块内存，使用__share__进行声明；
-
共享内存的大小为一个线程块中的线程数大小，其在使用的时候索引为当前线程的id，即threadIdx.x；
-
线程块中的线程同步使用***__syncthreads()函数，用于确保cache数组的写入操作在读取操作之前完成***。当一个线程使用__syncthreads()之后，能够确保线程块中的每个线程执行完__syncthreads()前面的语句后在执行后面的语句；
-
最后使用**归约（reduction）**算法将cache数组中的数相加；
-
polo
-
-
点积运算的优化
- 当某些线程需要执行一条指令，其他线程不需要执行时，称为线程发散；
- 如果将__syncthreads()放到if语句中，条件如果变为一部分能进入一部分线程不能进入，则会导致线程发散。这意味着有一部分线程无法完成__syncthreads()之前的操作，导致其永远无法执行__syncthreads()，从而进一步导致硬件是这些线程保持等待；
- 以上分析表明，将**__syncthreads()置于何处很关键**；
-
关于gridDim、blockDim、blockId、threadId等的关系
- gridDim.x和gridDim.y表示的是一个grid中的横向和纵向分别有有多少线程块block；
- blockDim.x和blockDim.y表示一个线程块block中的横向和纵向分别有多少线程；
- blockIdx.x和blockIdx.y表示该block处于横向和纵向的第几个；
- threadIdx.x和threadIdx.y表示线程位于block的第几行第几列；
-
polo
第六章 常量内存与事件
常量内存是GPU上的特殊内存，用于加速应用程序的执行；事件可以测试应用程序的性能，用于定量分析对应用程序的修改是否会带来性能的提升；
-
常量内存
-
cuda程序中除了支持全局内存和共享内存，还支持常量内存。其用于保存核函数执行期间不会发生变化的数据，其位于nvidia硬件中，大小为64KB；
-
光线追踪：从三维场景生成一张二维图像。原理为：在场景中的一个位置放置一台假想相机，其包含一个光传感器来生成图像，需要判断哪些光接触到该传感器。图像中的每个像素与命中传感器的光纤有着相同的颜色和强度，一般采用逆向计算，即假设从该像素发射的一道射线进入该场景，光线命中某个物体后计算该像素的颜色，即根据看到的颜色来设置颜色；
-
GPU上实现光线追踪：
-
代码
// CUDATest.cpp : 此文件包含 "main" 函数。程序执行将在此处开始并结束。 // #include <iostream> #include<cuda_runtime.h> #include<cuda_runtime_api.h> #include<device_launch_parameters.h> #include<device_functions.h> #include"cuda.h" #include"book.h" #include"cpu_bitmap.h" #define DIM 1024 #define rnd( x ) (x * rand() / RAND_MAX) #define INF 2e10f struct Sphere { float r, b, g; float radius; float x, y, z; __device__ float hit(float ox, float oy, float* n) { float dx = ox - x; float dy = oy - y; if (dx * dx + dy * dy < radius * radius) { float dz = sqrtf(radius * radius - dx * dx - dy * dy); //越大则表明越近 *n = dz / sqrtf(radius * radius); return dz + z; } return -INF; } }; #define SPHERES 20 __constant__ Sphere s[SPHERES]; //使用常量内存，直接指定所需内存大小 __global__ void kernel(unsigned char* ptr) { // map from threadIdx/BlockIdx to pixel position int x = threadIdx.x + blockIdx.x * blockDim.x; int y = threadIdx.y + blockIdx.y * blockDim.y; int offset = x + y * blockDim.x * gridDim.x; //计算偏移，gridDim.x表示横向有多少个线程块 float ox = (x - DIM / 2); float oy = (y - DIM / 2); float r = 0, g = 0, b = 0; float maxz = -INF; for (int i = 0; i < SPHERES; i++) { float n; float t = s[i].hit(ox, oy, &n); if (t > maxz) { float fscale = n; r = s[i].r * fscale; g = s[i].g * fscale; b = s[i].b * fscale; maxz = t; } } ptr[offset * 4 + 0] = (int)(r * 255); ptr[offset * 4 + 1] = (int)(g * 255); ptr[offset * 4 + 2] = (int)(b * 255); ptr[offset * 4 + 3] = 255; } // globals needed by the update routine struct DataBlock { unsigned char* dev_bitmap; }; int main(void) { DataBlock data; // capture the start time cudaEvent_t start, stop; //声明事件 HANDLE_ERROR(cudaEventCreate(&start)); HANDLE_ERROR(cudaEventCreate(&stop)); HANDLE_ERROR(cudaEventRecord(start, 0));//开始 CPUBitmap bitmap(DIM, DIM, &data); unsigned char* dev_bitmap; // allocate memory on the GPU for the output bitmap HANDLE_ERROR(cudaMalloc((void**)&dev_bitmap, bitmap.image_size())); // allocate temp memory, initialize it, copy to constant // memory on the GPU, then free our temp memory Sphere* temp_s = (Sphere*)malloc(sizeof(Sphere) * SPHERES); for (int i = 0; i < SPHERES; i++) { temp_s[i].r = rnd(1.0f); temp_s[i].g = rnd(1.0f); temp_s[i].b = rnd(1.0f); temp_s[i].x = rnd(1000.0f) - 500; temp_s[i].y = rnd(1000.0f) - 500; temp_s[i].z = rnd(1000.0f) - 500; temp_s[i].radius = rnd(100.0f) + 20; } HANDLE_ERROR(cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere) * SPHERES));//省去了给s分配和释放空间的代码 free(temp_s); // generate a bitmap from our sphere data dim3 grids(DIM / 16, DIM / 16); dim3 threads(16, 16); kernel << <grids, threads >> > (dev_bitmap); // copy our bitmap back from the GPU for display HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost)); // get stop time, and display the timing results HANDLE_ERROR(cudaEventRecord(stop, 0));//结束 HANDLE_ERROR(cudaEventSynchronize(stop)); float elapsedTime; HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop)); printf("Time to generate: %3.1f ms\n", elapsedTime); HANDLE_ERROR(cudaEventDestroy(start)); HANDLE_ERROR(cudaEventDestroy(stop)); HANDLE_ERROR(cudaFree(dev_bitmap)); // display bitmap.display_and_exit(); }
-
-
计算速度的吞吐量大于内存带宽的吞吐量，所以需要减少内存通信量；
-
-
常量内存用于保存在核函数执行期间不会发生变化的数据，Nvidia提供了64KB的常量内存，其不同于全局内存的处理，用其替换全局内存能够有效减少内存带宽；
-
为什么从常量内存中读取数据可以节约内存宽带？
- 对常量内存的单词读取操作可以广播到其他的“邻近”线程，节约了15的读取（参见线程束）；
- 常量内存的数据将缓存起来，相同地址的连续读操作不会产生额外的内存通信量；
-
线程束（warp）
- 在cuda架构中，线程束指的是一个包含32个线程的集合，这些线程以“步调一致”的形式执行；
- 线程束中的每个线程都将在不同的数据上执行相同的指令；
- 处理常量内存时，NVIDIA硬件将单次内存读取操作广播到每个半线程束，在半个线程束中包含了16个线程，半线程束中的每个线程都从常量内存的相同地址上读取数据，则GPU只会产生一次读取请求，并在随后将数据广播到每个线程；
- 需要注意的是，如果半线程束中的线程需要从不同的地址读取数据时，则会串行执行，降低效率，使得其低于使用全局内存；
- 从常量内存中读取大量的数据时，该方式的内存宽带为使用全局内存时的 1 / 16 1/16 1/16；
-
事件：用于分析程序性能
参考：《GPU高性能CUDA编程实战》
扫一扫
1203
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
