CUDA学习2-编程部分
CUDA编程
函数声明
host：主机端，通常指CPU
device：设备端，通常指GPU（数据可并行）
kernel：数据并行处理函数，在主机端调用kernel可以在设备端创建大量轻量级线程
host 和 device拥有各自的存储器
CUDA编程包括主机端和设备端两部分代码
| 执行位置 | 调用位置 | |
|---|---|---|
| _device_float DeviceFunc() | device | device |
| _global_void KernelFunc() | device | host |
| _host_float HostFunc() | host | host |
_global_定义一个kernel函数
入口函数，CPU上调用，GPU上执行
必须返回void
_device_and_host_可以同时使用
CUDA核函数（kernels）
在N个不同的CUDA线程上并行执行
//定义kernel_global_ void VecAdd(float* A, float* B, float* C){int i = threadIdx.x; C[i] = A[i] B[i];}int main(){//... //在N个不同的CUDA线程上并行执行 VecAdd<<<1, N>>>(A, B, C);}
线程层次（Thread Hierarchies）
块索引：blockIdx
维度：blockDim
一维、二维或三维
//单线程块//定义kernel_global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){int i = threadIdx.x; int j = threadIdx.y; C[i][j] = A[i][j] B[i][j];}int main(){//... //在N*N*1个不同的CUDA线程上并行执行 int numBlocks = 1; dim3 threadsPerBlock(N, N); VecAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);}//多线程块//定义kernel_global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){int i = threadIdx.x * blockDim.x threadIdx.x; int j = threadIdx.y * blockDim.y threadIdx.y; if(i<N && j<N) C[i][j] = A[i][j] B[i][j];}int main(){//... //并行执行 dim3 threadsPerBlock(16, 16); dim3 numBlocks(N/threadsPerBlock.x, N/threadsPerBlock.y); VecAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);}/*N = 32i = [0,1] * 16 [0,15]*/
CUDA内存传输
主机端可以从设备端往返传输数据
Global memory 全局存储器
Constant memory 常量存储器
cudaMalloc()：在设备端分配global memory
cudaFree()：释放存储空间
float *Md;//指向设备端上的一个存储空间int size = Width * Width * sizeof(float);cudaMalloc((void**)&Md, size);//...cudaFree(Md);
cudaMemcpy()：内存传输
cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);//参数：目的地址 源地址 大小 传输方向
host to host
host to device
device to host
device to device
例子：矩阵相乘
//CPU实现void MatrixMulOnHost(float* M, float* N, float* P, int width){for(int i=0; i<width; i) for(int j=0; j<width; j) {float sum = 0; for(int k=0; k<width; k) {float a = M[i * width k]; float b = N[k * width j]; sum = a*b; } p[i * width j] = sum; }}//cuda算法框架（3布）int main(void){//1.管理整个内存，为数据分配空间，将数据拷贝到GPU上 //2.在GPU上并行处理计算 //3.将结果拷贝回CPU}//GPU实现void MatrixMulOnDevice(float* M, float* N, float* P, int Width){int size = Width * Width * sizeof(float); //1.管理整个内存，为数据分配空间，将数据拷贝到GPU上 //分配输入 cudaMalloc(Md, size); cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); cudaMalloc(Nd, size); cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice); cudaMalloc(Pd, size); //2.在GPU上并行处理计算 _global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) {//访问一个matrix，采用二维block int tx = threadIdx.x; int ty = threadIdx.y; //每个kernel线程计算一个输出 float Pvalue = 0; //计算 for(int k=0; k<Width; k) {float Mdelement = Md[ty*Md.width k]; float Ndelement = Nd[k*Nd.width tx]; Pvalue = Mdelement Ndelement; } Pd[ty*Width tx] = Pvalue; } //3.将结果拷贝回CPU //1个block含width*width个线程 dim3 dimBlock(WIDTH, WIDTH); dim3 dimGrid(1, 1); MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd); //传送数据 cudaMemcpy(Pd, P, size, cudaMemcpyDeviceToHost); //释放 cudaFree(Md); cudaFree(Nd); cudaFree(Pd);}
主要性能问题：访存
赞 (0)