基本的cuda编程主要涉及两块
- 数据的传输处理
- Kernel函数的设计
GPU内部结构
粗浅来说,N卡中,GPU由多个SM处理器作为基本模块单元并包括GPU显存、L2外部缓存等。
每个SM处理器往往包括一个RT光追单元和四个子模块,每个子模块包括一个TensorCore和多个cudaCore(分为Int32、Fp32、Fp64)。
核函数的线程组织结构
Dim3类型的 Grid,子元素为Dim3类型的Block。每个Block包括若干个Thread。其中SM以一个线程束为单元进行运算,一个线程束包括32个Thread。为了提高计算率,Block应该是32的倍数。又一个Block只会在一个SM中进行运算,Block存在最大线程数限制,一般是1024。Grid的限制未知,一般很大,可以通过cuda系统API函数查询。
操作,通过 blockDim.x blockIdx.x等变量可以具体操作当前线程数据。
CUDA内存设计
首先CUDA的内存由定义在Host上的指针控制。
float* d_ptr
然后,主机中调用函数在cuda中分配内存
cudaMalloc ( void** devPtr, size_t size )
cudaMalloc( (void**)&d_ptr, size_t, N*sizeof(float) )
之后就是数据拷贝
cudaMemcpy( void* dst_ptr, void* src_ptr, size_t, cudaMemcpyKind)
cudaMemcpyAsync ( void* dst, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0 )
这里同步拷贝、和异步拷贝。更多的使用异步以充分利用带宽