CUDA- ZERO MEMORY: DEVICE ->HOST 메모리로 직접쓰기.

#include <stdio.h>
#include <stdlib.h>
#include <cutil_inline.h>
#include <cuda.h>

//벡터합 계산 커널
__global__ void vectorAdd(int *a, int *b, int *c)
{
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	c[idx] = a[idx] + b[idx];
}

int main()
{
	int i = 0;
	int nBlocks = 1024;
	int nThreads = 512;
	int Size = nBlocks* nThreads;
	size_t BufferSize = Size*sizeof(int);
	int *h_a, * h_b, * h_c;              // CPU 고정된 메모리 할당을 위한 포인터
	int *d_a, *d_b, *d_c;                // 메모리 맵핑을 위한 디바이스 포인터

	/* Allocate mapped CPU memory. */

	cudaHostAlloc((void **)& h_a, BufferSize, cudaHostAllocMapped);
	cudaHostAlloc((void **)& h_b, BufferSize, cudaHostAllocMapped);
	cudaHostAlloc((void **)& h_c, BufferSize, cudaHostAllocMapped);

	/* Initialize the vectors. */

	for(i = 0; i < Size; i++)
	{
		h_a[i] = i;
		h_b[i] = i;
	}

	//CPU 메모리를 맵핑하여 GPU 메모리 포인터로 가리키도록 한다.
	cudaHostGetDevicePointer((void **)&d_a, (void *)h_a, 0);
	cudaHostGetDevicePointer((void **)&d_b, (void *)h_b, 0);
	cudaHostGetDevicePointer((void **)&d_c, (void *)h_c, 0);

	////맵핑된 메모리 포인터를 이용하여 커널 호출

	vectorAdd<<< nBlocks, nThreads >>>(d_a, d_b, d_c);  
	cudaThreadSynchronize();

	cudaFreeHost(h_a);
	cudaFreeHost(h_b);
	cudaFreeHost(h_c);

	return 0;
}