奋斗了2周，终于把CUDA的内存与显存数据拷贝、pitch、以及如何对显存的数组进行引用弄明白了。很开心。
我是初次接触CUDA，学习CUDA是因为我了解到它的并行性使得数据量很大的程序的运行效率很高。我现在做的项目恰好需要这种高性能的并行运算。
我开始编写了几个内存和显存之间普通的int、char、float和double型数据的拷贝进行测试，然后进行了一维数组在内存和显存之间数据的拷贝，随后就是二维数组了。
在一维数组的定义和引用中，我开始是定义的时候指定数组的大小，后来改写为动态申请内存和显存。如下：
定义一个一维数组：
float *cpu_data; //内存中的一维数组指针
int width; //数组的宽度
cpu_data = (float*)malloc(sizeof(float) * width);
float *gpu_data; //显存中的一维数组指针
cutilSafeCall( cudaMalloc((void**) &gpu_data, sizeof(float) * width)); //为数组申请显存空间
for(int c = 0; c < width; ++c) //内存数组的初始化
cpu_data[c] = c;
//下面进行数据的拷贝 将内存数据拷贝到显存对应数组中
cutilSafeCall( cudaMemcpy( gpu_data, cpu_data, sizeof(float) * width, cudaMemcpyHostToDevice));
对于二维数组，在显存上申请空间是最好使用cudaMallocPitch（）函数，在此函数中有一个参数pitch，这个参数是补齐时显存数组每一行占得字节数。注意它是一个传出参数，其类型是size_t，也就是unsigned int类型。对于二维数组在内存和显存之间的拷贝最好使用函数
cudaMemcpy2D()，这样效率高些.
float *cpu_data; //CPU上的数组
float *gpu_data; //GPU上的数组
int width, height; //数组的宽度（列数）和高度（行数）
int pitch; //GPU上数组的pitch
printf("Input the width and height/n"); //输入数组的列数和行数
scanf("%d%d", &width, &height);
cpu_data = (float*)malloc(sizeof(float)*width*height); //申请内存空间
cutilSafeCall( cudaMallocPitch( (void**) &gpu_data, &pitch, sizeof(float) * width, height)); //申请显存空间
cutilSafeCall( cudaMemcpy2D( gpu_data, pitch, cpu_data, sizeof(float) * width, sizeof(float) * width, height, cudaMemcpyHostToDevice)); //内存和显存数据进行拷贝
下面是我写的最终的程序：
/************************************************************************** *矩阵相加的例子 ***************************************************************************/ #include <stdio.h> #include <stdlib.h> #include <cutil_inline.h>
/************************************************************************ * Init CUDA ************************************************************************/ #if __DEVICE_EMULATION__
bool InitCUDA(void){return true;}
#else bool InitCUDA(void) { int count = 0; int i = 0;
cudaGetDeviceCount(&count); if(count == 0) { fprintf(stderr, "There is no device./n"); return false; }
for(i = 0; i < count; i++) { cudaDeviceProp prop; if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) { if(prop.major >= 1) { break; } } } if(i == count) { fprintf(stderr, "There is no device supporting CUDA./n"); return false; } cudaSetDevice(i);
printf("CUDA initialized./n"); return true; }
#endif
/***************************************************************************************************** *kernel函数，矩阵相加 ******************************************************************************************************/ __global__ void myKernel(const float *a, const float *b, float *c, size_t pitch, int height, int width) { int i = blockIdx.y * blockDim.y + threadIdx.y; int j = blockIdx.x * blockDim.x + threadIdx.x; if(i < height && j < width) c[i * pitch/ sizeof(float) + j] = a[i * pitch / sizeof(float) + j] + b[i * pitch / sizeof(float) + j]; }
int main(int argc, char* argv[]) { if(!InitCUDA()) return 0; //CPU上的3个矩阵数组 float *cpu_A; float *cpu_B; float *cpu_C; //GPU上的3个矩阵数组 float *gpu_A; float *gpu_B; float *gpu_C; int width = 3; //矩阵的宽度（列数） int height = 2;//矩阵的高度（行数） size_t pitch; //GPU数组的pitch //为CPU上的矩阵数组申请内存空间 cpu_A = (float*)malloc(sizeof(float) * width * height); cpu_B = (float*)malloc(sizeof(float) * width * height); cpu_C = (float*)malloc(sizeof(float) * width * height); //为GPU上的矩阵数组申请显存空间 cutilSafeCall( cudaMallocPitch((void**) &gpu_A, &pitch, sizeof(float) * width, height)); cutilSafeCall( cudaMallocPitch((void**) &gpu_B, &pitch, sizeof(float) * width, height)); cutilSafeCall( cudaMallocPitch((void**) &gpu_C, &pitch, sizeof(float) * width, height)); //将pitch打印 printf("The pitch is: %d/n", pitch); //为CPU上的矩阵数组初始化 for(int r = 0; r < height; ++r){ for(int c = 0; c < width; ++c){ cpu_A[r * width + c] = r * c; cpu_B[r * width + c] = r + c; cpu_C[r * width + c] = 0.0; } } //打印CPU上的矩阵数组 printf("/nCPU_A DATA/n"); for(int r = 0; r < height; ++r){ for(int c = 0; c < width; ++c){ printf("%f/t", cpu_A[r * width + c]); } printf("/n"); } printf("/nCPU_B DATA/n"); for(int r = 0; r < height; ++r){ for(int c = 0; c < width; ++c){ printf("%f/t", cpu_B[r * width + c]); } printf("/n"); } printf("/nCPU_C DATA/n"); for(int r = 0; r < height; ++r){ for(int c = 0; c < width; ++c){ printf("%f/t", cpu_C[r * width + c]); } printf("/n"); }
//将CPU上的矩阵数组cpu_A、cpu_B分别拷贝到GPU上的矩阵数组gpu_A、gpu_B中 cutilSafeCall( cudaMemcpy2D( gpu_A, pitch, cpu_A, sizeof(float) * width, sizeof(float) * width, height, cudaMemcpyHostToDevice)); cutilSafeCall( cudaMemcpy2D( gpu_B, pitch, cpu_B, sizeof(float) * width, sizeof(float) * width, height, cudaMemcpyHostToDevice)); dim3 Dg(1, 2, 1); //定义整个grid的维度和尺寸 dim3 Db(width, 1, 1); //定义每个block的维度和尺寸 myKernel<<<Dg, Db, 0>>>(gpu_A, gpu_B, gpu_C, pitch, height, width); //调用kernel函数
//将显存数组gpu_C拷贝会内存数组cpu_C cutilSafeCall( cudaMemcpy2D( cpu_C, sizeof(float) * width, gpu_C, pitch, sizeof(float) * width, height, cudaMemcpyDeviceToHost));
//打印CPU_C数组 printf("/nAfter change CPU_C DATA/n"); for(int r = 0; r < height; ++r){ for(int c = 0; c < width; ++c){ printf("%f/t", cpu_C[r * width + c]); } printf("/n"); } //释放内存空间 free(cpu_A); free(cpu_B); free(cpu_C); //释放显存空间 cutilSafeCall( cudaFree(gpu_A)); cutilSafeCall( cudaFree(gpu_B)); cutilSafeCall( cudaFree(gpu_C)); //退出CUDA CUT_EXIT(argc, argv);
return 0; }
2周的编程（当然是闲暇时间，呵呵），我终于开始了我的CUDA编程之路。
路漫漫其修远兮，吾将上下而求索。
