CUDA Host<->Device cudaHostGetDevicePointer

대용량 데이터를 처리하다보니 GPU 메모리 제약으로 처리할 수 없어서 고민하던 중 Host메모리를 Device에서 사용 할 수 있는 방법을 찾았습니다. [cudaHostAlloc, cudaHostGetDevicePointer]

 

개발 환경은 Visual Studio 2013, CUDA ToolK 11.0, GPU 1050, Compute capability 6.1버전에서 테스트를 완료했습니다.

 

파일 2.5GByte를 Host Memory에 할당하고 CUDA Kernel에서 처리하였습니다. 

알고가면 좋은 내용을 간단히 정리하겠습니다. 1. Host에 할당한 메모리는 Read/Write가 모두 가능 2. Host Memory를 할당 시 공유 GPU 메모리 사용량 증가(Host 메모리의 절반 사용가능) 3. Device에서 Host 메모리를 사용하다보니 너무 잦은 메모리 I/O는 성능 저하!!

사용량을 확인하고 싶은 신 경우!!

빨간색 : cudaHostAlloc을 통해 할당 될 경우 공유 GPU 메모리 사용량이 증가합니다.

           [전체 메모리의 절반만큼 사용이 가능하고 Host메모리의 사용량이 증가합니다.]

파란색 : cudaMalloc은 전용 GPU 메모리 사용량!! -> GPU 메모리 제원만큼 사용이 가능!

보라색 : CUDA Kernel이 처리되면 사용량이 증가합니다!!

 

 

 

 

//////////////////////////////////////////
//////////////////////////////////////////
//! 파일을 바이너리에서 읽어옵니다.
BYTE* pbAlt = NULL;
__int64 lSize = 0;
FILE* fileAlt = _wfopen("파일명", L"r+b");
if (fileAlt != NULL)
{
	//read size of file
	_fseeki64(fileAlt, 0, SEEK_END);
	lSize = _ftelli64(fileAlt);
	_fseeki64(fileAlt, 0, SEEK_SET);

	pbAlt = new BYTE[lSize];

	unsigned int totnum = 0;
	unsigned int curnum = 0;
        
	//fread(&dAlt[0], sizeof(unsigned short), 1271012340, fileAlt);
	int nStepFile = 1024 * 64;
	while ((curnum = fread(&pbAlt[totnum], 1, nStepFile, fileAlt)) > 0) {
		totnum += curnum;
	}
	fclose(fileAlt);
}

//////////////////////////////////////////
//////////////////////////////////////////
//! CUDA관련

//! GPU Device상태
cudaDeviceProp m_deviceProp;

//! GPU가 한개 이기 때문에 0인덱스의 정보를 얻어서 HostMemory사용 여부를 확인합니다.
cudaGetDeviceProperties(&m_deviceProp, 0);

if (!m_deviceProp.canMapHostMemory)
{
	fprintf(stderr, "Device does not support mapping CPU host memory!\n");
	return false;
}


char *h_pchAlt = NULL;

if(pbAlt)
{
	//! Host메모리를 할당해 줍니다.(CUDA에서 쓰기위해서 필요합니다.)
	checkCudaErrors(cudaHostAlloc((void **)&h_pchAlt, lSize, cudaHostAllocWriteCombined));
    //! CUDA에서 할당해 준  Host 메모리에 복사해줍니다.
	memcpy(h_pchAlt, pbAlt, lSize);
	
    char *d_pchAlt;
	cudaHostGetDevicePointer((void **)&d_pchAlt, (void *)h_pchAlt, 0);
	
    //! 만든 커널을 통해 데이터를 사용하면 됩니다.(Read, Write 모두 가능)
    KernelTest << < 1, 1024 >> > (d_pchAlt);
    
    //! 커널이 종료될 때 까지 기다려줍니다.
    checkCudaErrors(cudaDeviceSynchronize());
}

cudaFreeHost(h_pchAlt);

delete[] pbAlt;
pbAlt = NULL;

 

c++에서 사용한 전체 로직이며 실제 사용한 소스는 아닙니다. 참고하면 좋을 것 같아서 첨부합니다. 
cudaHostGetDevicePointer 샘플소스는 "CUDA Samples\v11.0\1_Utilities\bandwidthTest"를 참고하는 것도 좋을 것 같습니다.