CUDA拷贝结构体成员指针

!!! 使用Unified Memory 神器避免这些 struct deep copy 把， 简单实用~~~

在网上看了好多CUDA中在拷贝结构体以及在内核中如何访问结构体成员变量的帖子，都没找到一个质量好的，so， 自己写一个吧。。。

先来个结构体：

struct Stu_Struct{
	int count;			//人数
	short* score;		//每个人的分数，总计count
	short num;			//附加分数
};

1. 说明

          首先大家都知道在主机端和设备端分别创建结构体指针，然后在主机端和设备端分配结构体成员指针的内存，最近从主机端拷贝结构体到设备端。这里先在主机端初始化一个结构体变量：

Stu_Struct h_stru;
	h_stru.count = 16;
	h_stru.score = (short*)malloc(sizeof(short) * h_stru.count);
	h_stru.num = 5;

接着，创建一个设备结构体指针，

    Stu_Struct* d_stru;
	cudaStatus = cudaMalloc((void**)&d_stru,  sizeof(Stu_Struct));
	if (cudaStatus != cudaSuccess) {
		printf("cudaMalloc failed!  Do you have a CUDA-capable GPU installed?");
	}

	cudaStatus = cudaMalloc((void**)&d_stru->score ,sizeof(short) * h_stru.count);
	if (cudaStatus != cudaSuccess) {
		printf("cudaMalloc failed!  Do you have a CUDA-capable GPU installed?");
	}

当然，这是一个错误的示范，，，包括很多网上博客也是这么乱写出来的。仔细看这段代码，首先是在设备端分配一个结构体变量和结构体成员score*，可是d_stru->score可以这么用吗？？如果你用调试会发现这是无效访问，显然在主机端试图访问设备端结构体指针成员变量score是不行的，d_stru是存在GPU上的，访问它的成员也只能在__global__函数中才可以的。但为了内核函数中可以使用结构体成员指针，又必须分配成员的内存，所以得在思路上绕一个弯。

      我们先分配一个d_score的设备指针，并将主机的 score的内容拷贝到设备d_score上：

	//分配设备结构体 成员指针
	short* d_score =NULL;
	cudaStatus = cudaMalloc((void**)&d_score, sizeof(short) * h_stru.count);
	CUDA_CHECK(cudaStatus);

	cudaStatus = cudaMemcpy(d_score, h_stru.score, sizeof(short) * h_stru.count, 
                             cudaMemcpyHostToDevice  );
	CUDA_CHECK(cudaStatus);

至此，我们在设备端分配了结构体d_stru和其 ‘成员’d_score的显存，但是二者并没有关联起来，我们需要将d_stru->score 与d_score关联其起来。这里需要一个指针进行中转，然后再复制结构体；

    // d_score 与 d_stru.score进行关联，需要一个临时指针进行中转
	short* temp = h_stru.score;
	h_stru.score = d_score;//****将主机指针指向设备指针  *****

	//结构体复制
	cudaStatus = cudaMemcpy(d_stru, &h_stru, sizeof(Stu_Struct), cudaMemcpyHostToDevice);
	CUDA_CHECK(cudaStatus);

	//恢复原指针
	h_stru.score = temp;

这样就可以在设备端为所欲为去访问结构体各个成员了。

   2.示例代码

完整示例程序请下载：

https://download.csdn.net/download/zhangfuliang123/11348369

*************************************************转载请注明出处,谢谢***********************************************