|
奋斗了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编程之路。
路漫漫其修远兮，吾将上下而求索。
| 
转播
