C++编程笔记(GPU并行编程-2)

C++与CUDA

内存管理

封装
利用标准库容器实现对GPU的内存管理

#include <iostream>
#include <cuda_runtime.h>
#include <vector>
#include <cstddef>
template<class T>
struct CUDA_Allocator {
  using value_type = T;  //分配器必须要有的
  T *allocate(size_t size) {
    T *dataPtr = nullptr;
    cudaError_t err = cudaMallocManaged(&dataPtr, size * sizeof(T));
    if (err != cudaSuccess) {
      return nullptr;
    }
    return dataPtr;
  }
  void deallocate(T *ptr, size_t size = 0) {
    cudaError_t err = cudaFree(ptr);
  }
};
__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 size = 65523;
  std::vector<int, CUDA_Allocator<int>> arr(size);
  kernel<<<13, 28>>>(arr.data(), size);
  cudaError_t err = cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  for (int i = 0; i < size; ++i) {
    printf("arr[%d]=%d\n", i, arr[i]);
  }
}

其中allocate和deallocate是必须实现的
这里不用默认的std::allocate,使用自己定义的分配器，使得内存分配在GPU上
vector是会自动初始化的，如果不想自动初始化的化，可以在分配器中自己写构造函数

官方提供的容器

#include <thrust/universal_vector.h>
int main(){
	//使用的是共享内存
  thrust::universal_vector<float> arr(size);
  }

或者

#include <thrust/host_vector.h>
#include <thrust/universal_vector.h>



  thrust::device_vector<float> dVec(100);
  //重载了=符号，会自动拷贝内存，这里是将GPU内存拷贝到CPU，
  thrust::host_vector<float> hVec = dVec;

关于分配器的更多介绍

函数调用

template<class Func>
__global__ void para_for(int n, Func func) {
  for (int i = blockDim.x * blockIdx.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x) {
    func(i);
  }
}
//定义一个仿函数
struct MyFunctor {
  __device__ void operator()(int i) {
    printf("number %d\n", i);
  }
};

int main() {
  int size = 65513;
  para_for<<<13,33>>>(size,MyFunctor{});
  cudaError_t err = cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
}

同样的，lambda也是被支持的，但是要先在cmake中开启

target_compile_options(${PROJECT_NAME} PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>)

lambda

lambda写法

  para_for<<<13, 33>>>(size, [] __device__(int i) { printf("number:%d\n", i); });

lambda捕获外部变量
一定要注意深拷贝和浅拷贝
如果这里直接捕获arr的话，是个深拷贝，这样是会出错的，因为拿到的arr是在CPU上的，而数据是在GPU上的，所以这里要浅拷贝指针，拿到指针的值，就是数据在GPU上的地址，这样就可以使用device函数对数据进行操作了

  std::vector<int, CUDA_Allocator<int>> arr(size);
  int*arr_ptr=arr.data();
  para_for<<<13, 33>>>(size, [=] __device__(int i) { arr_ptr[i] = i; });
  cudaError_t err = cudaDeviceSynchronize();
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
  for (int i = 0; i < size; ++i) {
    printf("arr[%d]=%d\n", i, arr[i]);
  }

同时还可以这样捕获

  para_for<<<13, 33>>>(size, [arr=arr.data()] __device__(int i) { arr[i] = i; });

时间测试

#include <chrono>
#define TICK(x) auto bench_##x = std::chrono::steady_clock::now();
#define TOCK(x) std::cout << #x ": " << std::chrono::duration_cast<std::chrono::duration<double> >(std::chrono::steady_clock::now() - bench_##x).count() << "s" << std::endl;

  
int main(){
  int size = 65513;

  std::vector<float, CUDA_Allocator<float>> arr(size);
  std::vector<float> cpu(size);

  TICK(cpu_sinf)
  for (int i = 0; i < size; ++i) {
    cpu[i] = sinf(i);
  }
  TOCK(cpu_sinf)

  TICK(gpu_sinf)
  para_for<<<16, 64>>>(
      size, [arr = arr.data()] __device__(int i) { arr[i] = sinf(i); });
  cudaError_t err = cudaDeviceSynchronize();
  TOCK(gpu_sinf)
  if (err != cudaSuccess) {
    printf("Error:%s\n", cudaGetErrorName(err));
    return 0;
  }
}

结果:

可以看到，求正弦GPU是要快于CPU的，这里差距还不明显，一般来说速度是由数量级上的差距的

学习链接