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是會自動初始化的，如果不想自動初始化的化，可以在分配器中自己寫建構函式
關於分配器的更多介紹

函式呼叫

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的，這裡差距還不明顯，一般來說速度是由數量級上的差距的

學習連結