C++编程笔记(GPU并行编程)
一、配置并使用
环境：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)
至此，就可以开始写代码了
二、代码
__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
}
一个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;
}
核函数可以嵌套，相互调用
__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;
}
三、内存管理
从核函数中返回数据
由于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();
}
如何解决无法获得程序的返回值这个问题呢？
指针？
#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;
}
好像不行
通过cudaGetErrorName获取到底为什么会出错
int ret = 0;
kernel<<<1, 1>>>(&ret);
cudaError_t err= cudaDeviceSynchronize();
printf("error : %s\n",cudaGetErrorName(err));
返回的是非法地址错误，类似于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;
}
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;
}
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;
比较新的显卡支持的特性，只将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;
}
使用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;
}
使用多线程对数组进行初始化
__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;
由于<<<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);
}
这样无论调用多少板块多少线程多少网格，都可以运行了
