C++與CUDA 內(nèi)存管理

封裝
利用標(biāo)準(zhǔn)庫容器實(shí)現(xiàn)對GPU的內(nèi)存管理

創(chuàng)新互聯(lián)公司主營邊壩網(wǎng)站建設(shè)的網(wǎng)絡(luò)公司,主營網(wǎng)站建設(shè)方案,成都app軟件開發(fā)公司,邊壩h5微信平臺(tái)小程序開發(fā)搭建,邊壩網(wǎng)站營銷推廣歡迎邊壩等地區(qū)企業(yè)咨詢

#include#include#include#includetemplatestruct 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>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是必須實(shí)現(xiàn)的
這里不用默認(rèn)的std::allocate,使用自己定義的分配器，使得內(nèi)存分配在GPU上
vector是會(huì)自動(dòng)初始化的，如果不想自動(dòng)初始化的化，可以在分配器中自己寫構(gòu)造函數(shù)
關(guān)于分配器的更多介紹

官方提供的容器

#includeint main(){//使用的是共享內(nèi)存
  thrust::universal_vectorarr(size);
  }

或者

#include#include  thrust::device_vectordVec(100);
  //重載了=符號(hào)，會(huì)自動(dòng)拷貝內(nèi)存，這里是將GPU內(nèi)存拷貝到CPU，
  thrust::host_vectorhVec = dVec;

函數(shù)調(diào)用

template__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);
  }
}
//定義一個(gè)仿函數(shù)
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 $<$:--extended-lambda>)

lambda

lambda寫法

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

lambda捕獲外部變量
一定要注意深拷貝和淺拷貝
如果這里直接捕獲arr的話，是個(gè)深拷貝，這樣是會(huì)出錯(cuò)的，因?yàn)槟玫降腶rr是在CPU上的，而數(shù)據(jù)是在GPU上的，所以這里要淺拷貝指針，拿到指針的值，就是數(shù)據(jù)在GPU上的地址，這樣就可以使用device函數(shù)對數(shù)據(jù)進(jìn)行操作了

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

同時(shí)還可以這樣捕獲

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

時(shí)間測試

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

  
int main(){int size = 65513;

  std::vector>arr(size);
  std::vectorcpu(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;
  }
}

結(jié)果:

可以看到，求正弦GPU是要快于CPU的，這里差距還不明顯，一般來說速度是由數(shù)量級(jí)上的差距的

學(xué)習(xí)鏈接

你是否還在尋找穩(wěn)定的海外服務(wù)器提供商？創(chuàng)新互聯(lián)www.cdcxhl.cn海外機(jī)房具備T級(jí)流量清洗系統(tǒng)配攻擊溯源，準(zhǔn)確流量調(diào)度確保服務(wù)器高可用性，企業(yè)級(jí)服務(wù)器適合批量采購，新人活動(dòng)首月15元起，快前往官網(wǎng)查看詳情吧

分享標(biāo)題：C++編程筆記(GPU并行編程-2)-創(chuàng)新互聯(lián)
地址分享：http://m.rwnh.cn/article0/copgoo.html

成都網(wǎng)站建設(shè)公司_創(chuàng)新互聯(lián)，為您提供品牌網(wǎng)站設(shè)計(jì)、網(wǎng)站改版、微信公眾號(hào)、網(wǎng)站內(nèi)鏈、網(wǎng)站設(shè)計(jì)公司、服務(wù)器托管

廣告

聲明：本網(wǎng)站發(fā)布的內(nèi)容（圖片、視頻和文字）以用戶投稿、用戶轉(zhuǎn)載內(nèi)容為主，如果涉及侵權(quán)請盡快告知，我們將會(huì)在第一時(shí)間刪除。文章觀點(diǎn)不代表本網(wǎng)站立場，如需處理請聯(lián)系客服。電話：028-86922220；郵箱：631063699@qq.com。內(nèi)容未經(jīng)允許不得轉(zhuǎn)載，或轉(zhuǎn)載時(shí)需注明來源： 創(chuàng)新互聯(lián)