C++编程笔记(GPU并行编程)_c++ gpu

内容来源：这个UP主课讲的很好小彭老师课程

一、配置并使用

环境：Windows10 + CLion + VS2019

cuda的安装，并行的话只需要安装cuda，cuDNN就不必了

编译器设置，windows下建议使用MSVC，因为是官方支持的，记得架构一定要设置amd64

GPU版本架构查询网址
CmakeList.txt编写

cmake_minimum_required(VERSION 3.22)#跟据自己的cmake版本来设置

project(CUDA_TEST2 LANGUAGES CXX CUDA)

set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_BUILD_TYPE Release)

set(CMAKE_CUDA_ARCHITECTURES 61)

#设置device函数的声名和定义分开,为了性能，不建议这样
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)

add_executable(CUDA_TEST2 main.cu)

#clion加的,好像可以不要
#set_target_properties(CUDA_TEST2 PROPERTIES
#        CUDA_SEPARABLE_COMPILATION ON)

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

至此，就可以开始写代码了

---

二、代码

• __global__：定义在GPU上的核函数，GPU执行，从CPU调用，可以有参数，不能有返回值，global就是GPU的main函数，device就是其他函数
• __device__：device修饰的函数定义在GPU上，device是设备函数，GPU执行，只能由GPU调用，调用规则与普通函数一样
• __host__：host修饰的函数定义在CPU上，只能在CPU上调用，调用规则与普通函数一样
• 同时加上host和device修饰符表示同时定义在CPU和GPU上，二者都可调用

__global__ void kernel() {
  //sayHi();
  printf("Block %d of %d, Thread %d of %d\n", blockIdx.x, gridDim.x, threadIdx.x, blockDim.x);
  //printf("Thread Numbers %d\n",blockDim.x);
  //线程编号
  //printf("Thread %d\n",threadIdx.x);
  //sayHello();
}

__device__ void sayHi() {
  printf("Hi , GPU\n");
}

__host__ void sayHiCpu() {
  printf("Hi , CPU\n");
}

__host__ __device__ void sayHello() {
  //从GPU上调用就有这个宏
#ifdef __CUDA_ARCH__
  printf("Hello , GPU CUDA_ARCH=%d\n", __CUDA_ARCH__);
#else
  printf("Hello , CPU NO CUDA_ARCH\n");
#endif
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

一个kernel就是一个网格
网格内有多个板块，板块内有多个线程组
核函数调用时，parm1:板块 parm2:线程数量

#include <iostream>
#include <cuda_runtime.h>

int main() {
  kernel<<<2, 3>>>();

  kernel<<<dim3(2, 2, 2), dim3(2, 2, 2)>>>();
  cudaDeviceSynchronize();

  return 0;
}
1
2
3
4
5
6
7
8
9
10
11

核函数可以嵌套，相互调用

__global__ void another() {
  printf("another : Thread %d of %d\n", threadIdx.x, blockDim.x);
}
__global__ void kernel_1() {
  printf("kernel_1 : thread %d of %d\n", threadIdx.x, blockDim.x);
  //可以在核函数里调用核函数
  another<<<1,3>>>();
}
int main(){
	 kernel_1<<<2, 3>>>();
	 return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12

三、内存管理

从核函数中返回数据
由于kernel的调用是异步的，我们创建一个kernel任务函数并调用后，它不会立刻返回，而是将任务提交到GPU的任务队列中，此时GPU并没有将任务执行完，和多线程有点像，所以不能直接从kernel函数中获取相应的返回值
所以核函数的返回类型必须为void
这样是会出错的，不被允许

__global__ int kernel() {
  return 100;
}
int main(){
  int ret = kernel<<<1, 1>>>();
  printf("ret=%d\n", ret);
  // kernel<<<dim3(2, 2, 2), dim3(2, 2, 2)>>>();
  cudaDeviceSynchronize();
  }
1
2
3
4
5
6
7
8
9

如何解决无法获得程序的返回值这个问题呢？
指针？

#include <iostream>
#include <cuda_runtime.h>
#include <common_functions.h>
__global__ void kernel(int *param) {
  *param = 100;
}

int main() {
  int ret = 0;
  kernel<<<1, 1>>>(&ret);
  cudaDeviceSynchronize();
  printf("ret=%d\n",ret);
  return 0;
}

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

好像不行

通过cudaGetErrorName获取到底为什么会出错

  int ret = 0;
  kernel<<<1, 1>>>(&ret);
 cudaError_t err= cudaDeviceSynchronize();
 printf("error : %s\n",cudaGetErrorName(err));
1
2
3
4

返回的是非法地址错误，类似于CPU上的段错误
因为这是堆栈上的局部变量，即使使用new或者malloc分配的内存，GPU也是访问不到的，因为GPU和CPU上的内存是独立的
GPU操作的其实是显存，当我们传一个地址过去的时候，GPU误以为是显存，而这块地址肯定是没有被分配过的

int main() {
  int *ret;
  kernel<<<1, 1>>>(cudaMalloc(&ret, sizeof(int)));
  cudaError_t err= cudaDeviceSynchronize();
  printf("error : %s\n",cudaGetErrorName(err));
  printf("ret=%d\n",ret);
  cudaFree(ret);
  return 0;
}
1
2
3
4
5
6
7
8
9

cudaMalloc可以实现申请GPU显存
可是CPU却访问不了显存的呀，所以这种方式也不行
解决办法

__global__ void kernel(int *param) {
  *param = 100;
}

int main() {
  int *ret;								//这里通过一个指针存放显存地址
  cudaMalloc(&ret,sizeof(int));			//将显存指针的地址传过去，传二级指针，使得指针可以指向一个显存地址
  kernel<<<1,1>>>(ret);
  cudaError_t err = cudaDeviceSynchronize();
  printf("error : %s\n", cudaGetErrorName(err));
  
  int retS;								//cpu上的内存
  //参数顺序是从右往左拷贝
  cudaMemcpy(&retS,ret, sizeof(int),cudaMemcpyDeviceToHost);//将GPU上的内容拷贝到CPU
  err = cudaDeviceSynchronize();
  printf("error : %s\n", cudaGetErrorName(err));
  printf("ret=%d\n", retS);
  cudaFree(ret);
  return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

cudaMemcpy会自动进行同步，所以上面代码的两个cudaDeviceSynchronize可以去掉
统一内存管理
使用cudaMallocManaged函数

  int *ret;
  cudaError_t err = cudaMallocManaged(&ret, sizeof(int));
  kernel<<<1, 1>>>(ret);
  cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  printf("ret=%d\n", *ret);
  cudaFree(ret);
  return 0;
1
2
3
4
5
6
7
8
9
10
11

比较新的显卡支持的特性，只将cudaMalloc换成cudaMallocManaged即可，这样分配出来的内存地址，无论在GPU和CPU上都是一样的，而且二者会同步变化，大大方便了开发人员。当然，这是有开销的

数组的分配

和平时的CPU分配数组差不多，只是需要替换成对应的API

__global__ void kernel(int *arr, int arrLen) {
  for (int i = 0; i < arrLen; ++i) {
    arr[i] = i;
  }
}
int main(){
  int *ret;
  int arrLen = 100;
  cudaError_t err = cudaMallocManaged(&ret, arrLen * sizeof(int));
  kernel<<<1, 1>>>(ret, arrLen);
  cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  for (int i = 0; i < arrLen; ++i)
    printf("arr[%d]=%d\n", i, ret[i]);
  cudaFree(ret);
  return 0;

}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21

使用cudaMalloc就应该是这样

int main() {
  int *ret;
  int arrLen = 100;
  //在GPU上申请数组内存
  cudaMalloc(&ret, arrLen * sizeof(int));
  //初始化数组
  kernel<<<1, 1>>>(ret, arrLen);
  cudaError_t err = cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  //在CPU上申请对应的内存并将GPU上的内容拷贝
  int *retS = (int *) malloc(arrLen * sizeof(int));
  err = cudaMemcpy(retS, ret, arrLen * sizeof(int), cudaMemcpyDeviceToHost);

  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  for (int i = 0; i < arrLen; ++i)
    printf("arr[%d]=%d\n", i, retS[i]);
  cudaFree(ret);
  free(retS);
  return 0;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

使用多线程对数组进行初始化

__global__ void kernel(int *arr, int arrLen) {
  int i = threadIdx.x;
  arr[i] = threadIdx.x;
}
int main() {
  int *ret;
  int arrLen = 100;
  cudaError_t err = cudaMallocManaged(&ret, arrLen * sizeof(int));
  //启动100个线程，给数组赋上线程序号
  kernel<<<1, arrLen>>>(ret, arrLen);
  cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  for (int i = 0; i < arrLen; ++i)
    printf("arr[%d]=%d\n", i, ret[i]);
  cudaFree(ret);
  return 0;
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19

由于<<<m,n>>>参数n太大可能会出错，因为不允许太多的线程，所以可以这样

__global__ void kernel(int *arr, int arrLen) {
//扁平化线程数量处理，网格跨步循环
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < arrLen; i += blockDim.x * gridDim.x) {
    arr[i] = i;
    printf("i=%d\n", i);
  }
}
int main(){
  int *ret;
  int N = 6553213;
  int blockNum = 312;
  int threadNum = (N + blockNum - 1) / N;	//向上取整
  cudaError_t err = cudaMallocManaged(&ret, N * sizeof(int));
  kernel<<<blockNum, threadNum>>>(ret, N);
  err = cudaDeviceSynchronize();
  for (int i = 0; i < N; ++i) {
    printf("arr[%d] = %d\n", i, ret[i]);
  }
  cudaFree(ret);
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20

这样无论调用多少板块多少线程多少网格，都可以运行了
学习链接