cuda数据传输之cudaMemcpy()和cudaMemcpy2D()详解

1. cudaMemcpy()函数传递一维数组方式：

num_elements = 1000;//1000个元素
double* host_a;//主机端的一维数组数据
double* host_result;//主机端的一维数组用于接收核函数计算结果
for(int i = 0; i<num_elements ; i++)
{
   host_a[i] = i; 
   host_result[i] = 0; 
}

//第一步：定义设备端一维指针用于接收主机内存数据
double* dev_a;
double* dev_result;

//第二步：定义一维数组所占内存字节大小
size_t size_element = num_elements * sizeof(double);//一个double占8个字节，共num_elements个元素

//第三步：分配设备端指针所指的显存空间大小：
cudaMallocManaged(&dev_a, size_element);
cudaMallocManaged(&host_result, size_element);

//第四步：将CPU内存数据传递给GPU显存数据
cudaStatus = cudaMemcpy(dev_a, host_a, size_element, cudaMemcpyHostToDevice);
cudaStatus = cudaMemcpy(dev_result, host_result, size_element, cudaMemcpyHostToDevice);

//定义一个block为256个线程，根据元素个数算出一个grid中分配多少个blocks，这里按一维分布。
dim3  blocks(256);
dim3  grid((num_elements_ + blocks.x - 1) / blocks.x);

// 选择哪个GPU执行工作。如果只有一个GPU，则肯定是零。
cudaStatus = cudaSetDevice(0);
kernel<<<blocks, grid>>>(dev_a, dev_result, num_elements);//执行核函数计算
cudaGetLastError();//检查核函数执行状况
cudaDeviceSynchronize();//同步操作

cudaMemcpy(host_result, dev_result, size_element, cudaMemcpyDeviceToHost);//计算结果传回

//释放设备段指针所申请的内存
cudaFree(dev_a);
cudaFree(dev_result);

//核函数实体
kernel<<<blocks, grid>>>(dev_a, num_elements);
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;//定位到当前线程
    if(i<num_elements)//num_elements个线程同时计算
    {
        dev_result[i] = dev_a[i];
    }
}

2. cudaMemcpy()函数传递二维数组方式：

num_elements = 10000;//100*100个元素
double** host_a= new double* [num_elements];//主机端的二维数组数据
	for (int i = 0; i< 100; i++)
	{
		host_a[i] = new double[100];
		for (int j = 0; j< 100;j++)
		{
			host_a[i][j] = i+j;
		}
	}
double* host_result;
	for (int i = 0; i< 10000; i++)
	{
		host_result[i] = 0;
	}

//第一步：定义设备端一维指针用于接收主机内存数据
double** dev_a;
double* dev_result;

//第二步：定义一维数组所占内存字节大小
size_t size_element = 100* sizeof(double*);//一个double*占8个字节，共100行

//第三步：分配设备端指针所指的显存空间大小：
cudaMallocManaged(&dev_a, size_element);
cudaMallocManaged(&host_result, 10000* sizeof(double));

//第四步：将CPU内存数据传递给GPU显存数据
	for (int i = 0; i < 100; i++)
	{
		size_t size_a = sizeof(double) * 100;//每一行所占的字节大小
		cudaMallocManaged(&dev_a[i], size_a);//给每一行分配显存空间
		cudaStatus = cudaMemcpy(dev_a[i], host_a[i], size_a, cudaMemcpyHostToDevice);//将每一行CPU数据传递给dev_a[i]指针所指的显存空间中
	}
cudaStatus = cudaMemcpy(dev_result, host_result, size_element, cudaMemcpyHostToDevice);

//定义一个block为256个线程，根据元素个数算出一个grid中分配多少个blocks，这里按一维分布。
dim3  blocks(256);
dim3  grid((num_elements_ + blocks.x - 1) / blocks.x);

// 选择哪个GPU执行工作。如果只有一个GPU，则肯定是零。
cudaStatus = cudaSetDevice(0);
kernel<<<blocks, grid>>>(dev_a, dev_result, num_elements);//执行核函数计算
cudaGetLastError();//检查核函数执行状况
cudaDeviceSynchronize();//同步操作

cudaMemcpy(host_result, dev_result, size_element, cudaMemcpyDeviceToHost);//计算结果传回

//释放设备段指针所申请的内存
cudaFree(dev_a);
cudaFree(dev_result);

//核函数实体
kernel<<<blocks, grid>>>(dev_a, num_elements);
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;//定位到当前线程
    if(i<100)//10000个线程同时计算
    {
        for(int j = 0;j<100;j++)
        {
           dev_result[i*100+j] = dev_a[i][j];
        }
    }
}

3. cudaMemcpy2D()传递一维数组、二维数组，以及核函数索引遍历的方法

double* host_a= new host_a[100];//一维数组数据
for (int i = 0; i < 100; i++)
{
	host_a[i] = 0.0;
}
double* dev_a;//定义设备段指针用于接收CPU数据



size_t width_a = 1 * sizeof(char);//定义数组宽度，一维数组，每行只有一个元素，宽度当然是 1 喽！
size_t height= 100;//一维数组共100个元素，即数组长度（或叫高度）
size_t pitch;  //这里直接声明pitch即可，后面配分显存空间的函数会自动算出来，一般是512字节，或者1024字节


//cudaMallocPitch分配显存空间，自动对齐内存，极大的加快了数据传输速度
cudaMallocPitch((void**)&dev_a, &pitch, width_a * sizeof(double), height);

//cudaMemcpy2D()来传递CUP数据到GPU
cudaMemcpy2D(dev_a, pitch, host_a, width_a * sizeof(double), width_a * sizeof(double), height, cudaMemcpyHostToDevice);

//上面是传递一维数组host_a


double* host_result;//用于接收计算结果，定义成了一维的数组来接收
for (int i = 0; i< 100; i++)
{
	host_result[i] = 0;
}
double* dev_result;//核函数中保存计算结果

//下面传递二维数组host_c

size_t width_c = 3 * sizeof(char);//定义100*3的二维数组，每行3个元素，所以宽度为3个字节
size_t height= 100;//一维数组共100个元素，即数组长度（或叫高度）
size_t pitch;
double** host_b= new double* [num_elements];//主机端的二维数组数据
for (int i = 0; i< 100; i++)
{
	host_b[i] = new double[100];
	for (int j = 0; j< 3;j++)
	{
		host_b[i][j] = i+j;
	}
}

//将二维数组host_b转化为一维数组host_c再传入设备端

double* host_c = new double[height* width_c ];//100*3个元素
for (int i = 0; i < height * width_c ; i++)//host_c初始化
{
	host_c[i] = 0.0;
}

for(int i = 0; i<height; i++ )
{
    for(j = 0; j < width_c ; j++)
    {
        host_c[i*width_c  + j] = host_b[i][j];
    ]
}

double* deb_c;//定义设备段指针 

//cudaMallocPitch分配显存空间，自动对其内存，极大的加快了数据传输速度
cudaMallocPitch((void**)&deb_c, &pitch, width_c  * sizeof(double), height);
cudaMallocPitch((void**)&dev_result, &pitch, 1* sizeof(double), height);

//cudaMemcpy2D()来传递CUP数据到GPU
cudaMemcpy2D(deb_c, pitch, host_c, width_c  * sizeof(double), width_c  * sizeof(double), height, cudaMemcpyHostToDevice);

cudaMemcpy2D(dev_result, pitch, host_result, 1* sizeof(double), 1* sizeof(double), height, cudaMemcpyHostToDevice);

dim3  blocks(32);
dim3 grid((num_elements_ + blocks.x - 1) / blocks.x);
kernel<<<blocks, grid>>>(dev_a, width_a, width_c, pitch, height);

cudaGetLastError();
cudaDeviceSynchronize();

cudaMemcpy2D(host_result, width_a * sizeof(double), dev_result, pitch, width_a * sizeof(double), height, cudaMemcpyDeviceToHost);


//核函数实体
kernel<<<blocks, grid>>>(dev_a, deb_c, dev_result，width_a, width_c, pitch, height);
{

    double* a;
    double* c;
    double* result;

    int i = blockIdx.x * blockDim.x + threadIdx.x;//定位到当前线程
    if(i<100)//个线程同时计算
    {
        a = (double*)((char*)dev_a + i*pitch);//a指针指向每行存储的空间，每行有1个元素
        c = (double*)((char*)dev_c + i*pitch);//c指针指向每行存储的空间，每行有3个元素
        result = (double*)((char*)dev_result+ i*pitch);//result指针指向每行存储的空间，每行有3个元素
        for(int j = 0;j<3;j++)
        {
           result[0] += a[0] + c[j];//因为dev_a、dev_result是一维数组，每行只有一个元素，所以索引下标直接是0，而不需要像c[j]一样来遍历
//再者计算的result[0]直接改变了dev_result的对应的值，不需要dev_result[i] = result[0]，赋值操作，因为result是指针，直接指向dev_result内存，并对其进行计算。
        }
    }
}

总结：无论一维还是二维数组，都变为一维数组，进行传递，二维数组要注意数组的宽度大小。

核函数中：a = (double*)((char*)dev_a + i*pitch)，这样对a的遍历方式是获取数组每一行的数据，对一维数组每行的索引，直接是0下标索引，同时计算a的结果直接作用在dev_a 上，而不需要再将a赋值给dev_a！！！！