Stream

CUDA 프로그램의 구조

1. 호스트에서 디바이스로 입력 데이터 전송

2. 커널 함수에서 데이터 처리

3. 결과 데이터를 디바이스에서 호스트로 전송

 

위 과정을 차례대로 처리하면 입력 데이터가 전송되는 시간동안 GPU 프로세서는 대기하게 된다. 만일 입력 데이터를 좀 더 작게 나누고 입력 데이터 전송이 완료된 것부터 GPU에서 계산을 하는 동시에 데이터를 전송하면 프로그램의 전체 효율이 좋아진다. 데이터 입력을 받으면서 GPU 프로세서가 함께 계산하며 동시성을 높여 주는 것이 CUDA의 스트림이다. 

 

데이터 처리 과정을 여러 개로 나누어 데이터 입력 과정과 동시에 처리할 수 있도록 하려면 데이터의 의존성(Dependency)이 없어야 한다. 앞서 입력된 데이터의 결과가 뒤에 입력될 데이터에 사용되지 않아야 독립적으로 동시에 처리할 수 있다. CUDA 스트림은 의존 관계가 있는 데이터 처리에서도 실행 순서를 지킬 수 있는 기능도 제공하여 편리하게 사용할 수 있다. 

 

병행 복사  

데이터 복사와 실행을 병렬 처리하려면 복사하는 데이터의 크기를 작게 나누고, 비동기 복사 함수를 호출해야 한다.

//Asynchronous copy 
cudaError_t cudaMemcpyAsync( void * dst, const void * src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream = 0)

cudaMemcpyAsync()를 사용하는 방법은 cudaMemcpy()와 비슷하지만 4번째 인자로 스트림 객체를 넣을 수 있고 전송 시 사용하는 호스트 메모리는 cudaMallocHost()함수를 이용하여 생성한 고정 메모리를 사용해야 한다. 작게 나눈 데이터를 디바이스에 복사하고서 커널 함수가 실행되어 해당하는 데이터를 계산한다. 복사가 완료되기 전에 커널 함수가 실행되면 프로그램의 오류가 발생하게 되고 복사가 완료되고 나서도 커널 함수가 실행되지 않으면 병행 실행의 효과가 떨어진다. 

 이 순서를 맞추는 기능을 하는 것이 스트림 객체이다. 스트림 객체는 cudaStreamCreate()함수로 생성하고 cudaStreamDestroy()함수로 소멸시킨다. 

//Stream 객체 생성 및 소멸 
cudaError_t cudaStreamCreate( cudaStream_t *pStream);
cudaError_t cudaStreamDestroy( cudaStream_t stream);

 

//스트림 객체 생성, 소멸
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

cudaMemcpyAsync(deviceInput1, hostInput1, dataSize1, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(deviceInput2, hostInput2, dataSize2, cudaMemcpyHostToDevice, stream2);

kernel<<<nBlocks, nThreads, 0, stream1>>>(deviceInput1, deviceOutput1);
kernel<<<nBlocks, nThreads, 0, stream2>>>(deviceInput2, deviceOutput2);

cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);