关于CUDA中GPU内存的创建和传递

 1. cudaMalloc 和 cudaMemcpy

(1) cudaMalloc

cudaMalloc	(	void **	devPtr,
		size_t	size
	)

第一个参数是，指针的指针，感觉可以形成自己的范式。如：

        先定义一个指针 int* d_pit，如果想要让这个指针指向 device 上的内存分配地址，那么就可以按照下面这种方式调用 cudaMalloc 函数。

cudaMalloc( &d_pit, sizeof(int));

这里稍微回忆一下指针的两个运算符：* 和 &：

        假设我们创建了一个指针 a ，那么 *a 就是 a 所指向的地址上存放的值； &a 是 a 这个指针本身存储的地址。

(2) cudaMemcpy

cudaMemcpy	(	void *	dst,
		const void *	src,
		size_t	count,
		enum cudaMemcpyKind	kind
	)

前两个参数都是指针，数据传递的方向是从第二个参数指向的存储位置传递给第一个参数指向的存储位置，应和 kind 相对应。kind 有四种：

        cudaMemcpyHostToHost

        cudaMemcpyHostToDevice

        cudaMemcpyDeviceToHost

        cudaMemcpyDeviceToDevice

示例代码如下，很容易看懂，有很多库不用调，但是懒得删了......

#include "cuComplex.h"
#include <cufft.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include<math.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <iostream>
#include <vector>
#include <cmath>
#include <iomanip>

//控制参数

__global__ void yanzheng(int* d_pit)
{
	printf(" %d ", *d_pit);
}

int main()
{
	//Parameters on device
	int* d_pit;
	int* pit;
	pit = (int*)malloc(sizeof(int));	
	*pit = 1;

	cudaMalloc( &d_pit, sizeof(int));
	cudaMemcpy(d_pit, pit, sizeof(int), cudaMemcpyHostToDevice);
	yanzheng <<< 1, 1 >>>(d_pit);


	return 0;
}

 2. cudaMallocPitch 和 cudaMemcpy2D

        在处理二维数组时，我们可以借助这两个函数来实现数据的自动对齐，从而减少数据访问的时间。这两个函数的使用方法和 1 中的那对基本一样。多了 pitch 这个参数，这个参数存储的是，在设备上一行数据所占用的字节数。

官方文档如下：

cudaError_t cudaMallocPitch	(	void **	devPtr,
		size_t *	pitch,
		size_t	width,
		size_t	height
	)

cudaError_t cudaMemcpy2D	(	void *	dst,
		size_t	dpitch,
		const void *	src,
		size_t	spitch,
		size_t	width,
		size_t	height,
		enum cudaMemcpyKind	kind
	)

cudaMallocPitch 函数实现在 GPU 上分配内存，你指定了二维数组的一行的字节数和行数（width和height），给这个函数传递一个用于接受地址的指针和一个 size_t 类型的地址 &pitch，这个函数便会返回一个地址和分配结果（设备上一行数据所占用的字节数，存储在pitch）。

cudaMemcpy2D 实现数据的传递，从 src 传到 dst，需要告诉这个函数：你已经得到的 pitch ，以及 主机端的数组 一行的字节数，要传递的一行的字节数和行数。主机端的数组一行的字节数和要穿传递的一行的字节数往往是相等的。

目前我所掌握的方法是将一维数组传递到GPU的一维数组中（因此应将二维数组按行主序存在一个一维数组中，然后再将这个一维数组传递到GPU里，注意是按行主序存储哦）。

当正确传递之后，一个比较关键的点是如何实现数据的索引：

double* hhhRow = (double*) ((char*) hhh_2d+ (i) * pitch )

上面这行代码实现了如何找到第 i 行开头的地址，

然后 hhhRow[j],便可以索引到 第 i 行，第 j 列的数据。

作为例子，下面是实现求两个矩阵各个元素平方和后的矩阵的代码：

将在工作文件夹生成一个 txt 文件。

#include "cuComplex.h"
#include <cufft.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include<math.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <iostream>
#include <vector>
#include <cmath>
#include <iomanip>

#define N 8
int n_threadx = 1, n_thready = 1;
int n_blockx = N, n_blocky = N;


__global__ void square_kernel(double* Kx, double* Ky, double* Ksq, size_t pitch)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int j = blockIdx.y * blockDim.y + threadIdx.y;

	if (i < N && j < N)
	{	
		double* KsqRow = (double*)((char*)Ksq + (i)*pitch);
		double* KxRow = (double*)((char*)Kx + (i)*pitch);
		double* KyRow = (double*)((char*)Ky + (i)*pitch);
		KsqRow[j] = KxRow[j] * KxRow[j] + KyRow[j] * KyRow[j];
	}
}

//dim3
dim3 dimblock(n_threadx, n_threadx);
dim3 dimgrid((N + dimblock.x - 1) / dimblock.x, (N + dimblock.y - 1) / dimblock.y);


int main()
{
	//Parameters on host
	int h_NX = N;
	size_t pitch;
	double * X, * Y, * XYsq;

	X = new double[h_NX * h_NX];
	Y = new double[h_NX * h_NX];
	XYsq = new double[h_NX * h_NX];

	//Parameters on device
	double* d_X;
	double* d_Y;
	double* d_XYsq;
	size_t* d_pit;

	// compute X & Y : grid of coordinate_x & coordinate_y, 一定是行主序！
	for (int i = 0; i < h_NX; ++i) {
		for (int j = 0; j < h_NX; ++j) {
			X[j + i * h_NX] = 1.0 * i;
		}
	}

	for (int i = 0; i < h_NX; ++i) {
		for (int j = 0; j < h_NX; ++j) {
			Y[j + i * h_NX] = j * 1.0;
		}
	}


	//在设备上分配内存并数据传输
	cudaMallocPitch((void**)&d_X, &pitch, sizeof(double) * h_NX, h_NX);
	cudaMemcpy2D(d_X, pitch, X, sizeof(double) * h_NX, sizeof(double) * h_NX, h_NX, cudaMemcpyHostToDevice);
	cudaMallocPitch((void**)&d_Y, &pitch, sizeof(double) * h_NX, h_NX);
	cudaMemcpy2D(d_Y, pitch, Y, sizeof(double) * h_NX, sizeof(double) * h_NX, h_NX, cudaMemcpyHostToDevice);
	cudaMallocPitch((void**)&d_XYsq, &pitch, sizeof(double) * h_NX, h_NX);
	cudaMemcpy2D(d_XYsq, pitch, XYsq, sizeof(double) * h_NX, sizeof(double) * h_NX, h_NX, cudaMemcpyHostToDevice);

	//调用核函数
	square_kernel << < dimgrid, dimblock >> > (d_X, d_Y, d_XYsq, pitch);
	cudaDeviceSynchronize();
	cudaMemcpy2D( XYsq, sizeof(double) * h_NX, d_XYsq, pitch, sizeof(double) * h_NX, h_NX, cudaMemcpyDeviceToHost);

	FILE* fp;
	fp = fopen("pit.txt", "w");
	for (int i = 0; i < h_NX; i++)
	{
		for (int j = 0; j < h_NX; j++)
		{
			fprintf(fp, "%f  ", XYsq[j + i * h_NX]);

		}
		fprintf(fp, "\n");
	}
	fclose(fp);

	cudaFree(d_X);
	cudaFree(d_Y);
	cudaFree(d_XYsq);
	delete(X, Y, XYsq);

	return 0;
}