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

寫法

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]);
  }
           

同時還可以這樣捕獲

時間測試

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

學習連結