cuda 为包含指针的结构数组指针分配统一内存(UM)

使用统一内存后，程序会简化很多，不需要针对同一组数据定义两个数组（host and device），而且不需要显式地进行主机与设备之间的数据传输。至于程序的性能，很难通过简单的程序进行测试。如果只是针对核函数来说，可以说两个版本的核函数具有同样的性能。此外，统一内存的某些功能在Windows操作系统下依然受到限制。

将包含指针的结构体从host传进device，需要deep copy，十分麻烦，如果一开始就选择创建统一内存下的结构体，就不再需要deep copy。

需要分配两次，一是为结构数组指针分配，二是为结构数组包含的指针再分配一次，切记。

struct TEST {
	float* x;
	size_t num;
};

__global__ void testKernel(TEST* t) {
	for (int i = 0; i < t->num; i++) printf("%10.3f\n", t->x[i]);
}

int main()
{
	TEST* t;
	CHECK(cudaMallocManaged(&t, sizeof(TEST))); // 1. allocate UM for t
	t->num = 10;
	CHECK(cudaMallocManaged(&(t->x), sizeof(float) * t->num)); // 2. allocate UM for t->x
	
	float* y = new float[10];
	for (int i = 0; i < 10; i++) {
		y[i] = i;
	}
	memcpy(t->x, y, sizeof(float) * t->num);

	testKernel << <1, 1 >> > (t);
	CHECKKERNEL;
	CHECK(cudaDeviceSynchronize());

	return 0;
}

不要写成 TEST* t = new TEST; 传进kernel后会出现cuda error700。

全部使用统一内存进行debug非常方便，把__global__去掉，直接可以串行debug kernel函数。