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的,这里差距还不明显,一般来说速度是由数量级上的差距的
【推荐】编程新体验,更懂你的AI,立即体验豆包MarsCode编程助手
【推荐】凌霞软件回馈社区,博客园 & 1Panel & Halo 联合会员上线
【推荐】抖音旗下AI助手豆包,你的智能百科全书,全免费不限次数
【推荐】博客园社区专享云产品让利特惠,阿里云新客6.5折上折
【推荐】轻量又高性能的 SSH 工具 IShell:AI 加持,快人一步
· 在鹅厂做java开发是什么体验
· 百万级群聊的设计实践
· WPF到Web的无缝过渡:英雄联盟客户端的OpenSilver迁移实战
· 永远不要相信用户的输入:从 SQL 注入攻防看输入验证的重要性
· 全网最简单!3分钟用满血DeepSeek R1开发一款AI智能客服,零代码轻松接入微信、公众号、小程