文章目录

• • • 第一章
    • 第二章
    • 第三章
    • • 3.1 and 3.2
      • 3.3 查询设备
    • 第四章
    • 第五章 线程协作
    • 第六章 常量内存与事件

第一章

• 用途：加速三维成像，计算流体力学，分子模拟计算等等，这些应用场景大都需要进行大量的科学计算，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上实现光线追踪：
    • …
  • 计算速度的吞吐量大于内存带宽的吞吐量，所以需要减少内存通信量；

• 常量内存用于保存在核函数执行期间不会发生变化的数据，Nvidia提供了64KB的常量内存，其不同于全局内存的处理，用其替换全局内存能够有效减少内存带宽；

• 事件：用于分析程序性能
  参考：《GPU高性能CUDA编程实战》