GPU高性能编程CUDA入门

CMake中启用CUDA支持

最新版的 CMake(3.18 以上)，只需在LANGUAGES 后面加上 CUDA 即可启用。
然后在 add executable 里直接加你的 .cu文件，和 .cpp 一样。

cmake minimum required(VERSION 3.10)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE BUILD TYPE Release)

project(hellocuda LANGUAGES CXX CUDA)

add executable(main main.cu)

CUDA和C++的关系就像C++和C的关系一样，大部分都兼容，因此能很方便的重用C++现有的任何代码库，引用C++头文件等

第一个程序

将CPU以及系统的内存成为主机（host）。而将GPU及其内存成为设备（device）。
在GPU设备上执行的函数通常称为核函数（kernel）。
没有核函数，只考虑在主机运行的CUDA代码和标准的C在很大程度上是没有区别的。

#include <iostream>

__global__ void kernel(void){

}

int main(void){
  kernel<<<1,3>>>();
  printf("Hello,World!\n");
  return 0;
}

而当遇到具有__global__修饰符的函数时，编译器就会将该函数编译为在device上运行。在此例子中，函数kernel()将被交给编译device代码的编译器，而main()函数将被交给host编译器。
__inline__ 内联函数
__forceinline__ 强制内联
__noinline__ 禁止内联优化
__host__ 定义支持CPU调用
__device__ 定义支持GPU调用, 可以与__host__ 同时修饰

而对kernel()函数的调用语句则使用了一种尖括号和两个数值的方式，这里的<<<1, 3>>>

• 参数一：表示设备在执行款核函数时使用的并行线程块的数量。
• 参数二：需要多少个线程格（Grid）（一格表示N个线程块的集合）

参数传递

以下代码展示了如何像核函数传递参数并取得返回结果

#include "stdio.h"

__global__ void add(int a, int b, int *c){
  *c = a+b;
}

int main(){
  int c;
  int *dev_c;
  HANDLE_ERROR(cudaMalloc((void**)&dev_c, size(int)));

  add<<<1,1>>>(2,7,dev_c);

  HANDLE_ERROR(cudaMemcpy(&c,
    dev_c,
    sizeof(int),
    cudaMemcpyDeviceToHost));
   printf("2+7 = %d\n",c);
   cudaFree(dev_c);
   return 0;
}

上述代码说明，cuda可以像调用C函数那样将参数传递给核函数；当设备执行任何有用的操作时，都需要分配内存，例如将计算值返回给主机

cudaMalloc()函数：（注意，分配内存的指针不是该函数的返回值，这点与malloc()不同）

• 参数一： 一个指针，指向用于保存新分配内存地址的变量。注意，由于C语言中，指针传递是本身也是值传递的，所以为了使指针本身的值（不是指针地址指向的值）可以改变，因此在传递时要使用双重指针void**，这样做的主要原因还是因为分配内存的指针最终不是通过函数返回，而是直接改变参数值导致的（如果传的是一重指针，则改变的是pd指向的内存空间的数据，而不是pd本身，所以pd也就不能指向GPU的内存了）。
• 参数二：分配内存的大小

CUDA中对设备指针的使用限制总结如下：

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

在主机代码中，可以通过调用cudaMemcpy()来访问设备上的内存。这个函数调用的行为类型与标准C中的memcpy()，只不过多了一个参数来指定设备内存指针究竟是源指针还是目标指针。如，当最后一个参数为cudaMemcpyDeviceToHost时，代表运行时源指针是一个设备指针，而目标指针是以个主机指针。此外还有参数cudaMemcpyHostToDevice和cudaMemcpyDeviceToDevice等，如果源指针和目标指针都是位于主机上，那么可以直接调用标准C的memcpy()函数。

查询设备

对于拥有多个支持CUDA的设备，需要通过某种方式来确定使用的是哪一个设备。

int count;
HANDLE_ERROR(cudaGetDeviceCount(&count));

在调用cudaGetDeviceCount()后，可以对每个设备进行迭代，并查询各个设备的相关信息。CUDA runtime将返回一个cudaDeviceProp类型的结构，其中包含了设备的相关属性。可以利用cudaGetDeviceProperties()来获得i号设备的属性:

#include <iostream>

int main(){
    cudaDeviceProp prop;
    int count;
    cudaGetDeviceCount(&count);
    for(int i =0 ; i< count; i++){
	cudaGetDeviceProperties(&prop, i);
  //对设备的属性执行某些操作
    }
    std::cout<<count<<std::endl;
}

在知道了每个可用的属性以后，接下来就可以进行一些具体的操作，如：

std::cout<<prop.major<<std::endl;

设备属性的使用

根据在cudaGetDeviceCount()和cudaGetDeviceProperties()中返回的结果，我们可以对每个设备进行迭代，来找到我们期望的某些达到要求的设备。但是这种迭代操作执行起来有些繁琐，因此CUDA runtime提供了一种自动方式来执行这个迭代操作。首先，找出希望设备拥有的属性并将这些属性填充到一个cudaDeviceProp结构。

cudaDeviceProp prop;
memset(&prop, 0 , sizeof(cudaDeviceProp));
prop.major = 1;
prop.minor = 3;

之后，将该结构传递给cudaChooseDevice()，这样CUDA runtime运行时将查找是否存在某个设备满足这些条件，并返回一个设备ID，我们可以将这个设备ID传递给cudaSetDevice()。随后，所有的设备操作都将在这个设备上执行。

#include <iostream>
using std::cout;
using std::endl;

int main(){
  cudaDeviceProp prop;
  memset(&prop, 0 , sizeof(cudaDeviceProp));
  prop.major = 1;
  prop.minor = 3;

  int dev;
  cudaGetDevice(&dev);

  cudaChooseDevice(&dev, &prop);
  cout<<"ID:"<<dev<<endl;
}

CUDA C并行编程

基于GPU的矢量求和

#define N 10
int main(){
  int a[N],b[N],c[N];
  int *dev_a, *dev_b, *dev_c;

  //在GPU上分配内存，注意这里要知道为什么使用void**
  cudaMalloc( (void**)&dev_a, N*sizeof(int));
  cudaMalloc( (void**)&dev_a, N*sizeof(int));
  cudaMalloc( (void**)&dev_a, N*sizeof(int));

  ...//创建a，b数组并赋值

  //将数组a，b复制到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);

  //将数组c从GPU复制到CPU
  cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost);

  ...//显式结果

  //释放GPU上分配的内存
  cudaFree(dev_a);
  cudaFree(dev_b);
  cudaFree(dev_c);

  return 0;
}

看一下核函数的调用：
add<<<N,1>>>(dev_a, dev_b, dev_c);
尖括号中的两个数值将传递给runtime，作用是告诉runtime如何启动核函数：我们将每个并行执行环境都称为一个线程块（Block），对于此例，将有N个线程块在GPU上运行（N个运行核函数的副本）。

如何在代码中知道当前正在运行的是哪一个线程块？答：利用变量blockIdx.x
threadIdx.x 读取线程
blockDim.x 获取当前线程数量

add()函数：

__global__ void add(int *a, int *b, int *c){
  int tid = blockIdx.x;  //计算机该索引处的数据
  if(tid < N)
  c[tid] = a[tid] + b[tid];
}

当启动核函数时，我们将并行线程块的数量指定为N。这个并行线程块集合就称为一个“线程格（Grid）”。因此，此例表示我们想要一个一维的线程格，其中每个线程格包含N个线程块，每个线程块的blockInx.x的值都是不同的，cuda会为每个设备代码副本提供不同的blockInx.x。

需要注意的一点是：在启动线程块数组时，数组每一维（N）的最大数量不能超过65535。这是一种硬件限制，如过启动的线程块数量超过了这个限制，那么程序将运行失败。

线程协作

【并行计算】CUDA在现代C++中如何运用？看这一个就够了！
《GPU高性能编程CUDA实战 CUDA By Example》