Yes! Programmer

커널함수는 비동기 함수이기 때문에, 동기화 처리시, cudaDeviceSynchronie() 함수를 호출하면 된다.

근거 없는 자신감~
전엔 참 부정적이였는데....

근거 없는 자신감이 새로운 것을 창조하기 위한 근간이 된다는 것을 느끼는 순간이다.

 - 소개.

   쿠다는 아래와 같이 처리를 하는데, 데이터를 복사하는 과정과 GPU 처리 과정이 순차적으로 처리가 되기 때문에, 1번 과정에서 GPU는 대기하게 되어, 동시성이 떨어진다.

   이와 같이  대기 시간을 줄이기 위해, 데이터를 작은 단위로 나뉘고 데이터 전송이 완료된 것 부터 GPU에서 계산하는 동시에 계속해서 데이터를 전송하도록 개선 했는데,  이를 스트림이라고 한다.

   1. 호스트에서 디바이스로 입력 데이터 복사.

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

   3. 처리결과를 디바이스->호스트로 복사.

 -특징

    -  순차적으로 처리를 해야 하는 의존적인 데이터에 대해서도 스트림은 순차적으로 처리를 한다.

    -  cudaMallocHost() 를 이용해서 고정된 메모리[pinned Memory] 사용해야 한다.

    -  커널함수<<<1,2,3,stream>>>() - 4번재 항목에 적용된다. 명시하지 않으면 기본값 stream 0 이 할당된다.

- 관련 API

   1. cudaStreamCreate()   

   2. cudaMemcpyAsync();

   3. cudaStreamDestroy()

- 샘플코드

__global__ void kernel( int*In, int*Out)
{
    int tid = blockIdx.x * blockDim.x+ threadIdx.x;

    for(int i=0;i<5;i++)
        Out[tid] += In[tid];
}

int main()
{
    const int nStreams = 15;        //스트림 분할 개수
    const int nBlocks = 65535;      //블록의 개수
    const int nThreads = 512;       //스레드의 개수
    const int N = 512*65535;        //데이터의 개수
    const int Size = N*sizeof(int); //버퍼의 사이즈

    int* host_In;
    int* host_Out;

    //호스트 메모리 할당
    cudaMallocHost((void**)&host_In,Size);
    cudaMallocHost((void**)&host_Out,Size);

    //데이터 입력
    for( int i = 0; i < N; i++)
    {
        host_In[i] = i;
        host_Out[i] = 0;
    }

    int* dev_In;
    int* dev_Out;

    //디바이스 메모리 할당
    cudaMalloc((void**)&dev_In, Size);
    cudaMalloc((void**)&dev_Out, Size);

    cudaMemset(dev_In, 0, Size);
    cudaMemset(dev_Out, 0, Size);
       
    
    //스트림 객체 생성
    cudaStream_t *streams = (cudaStream_t*) malloc(nStreams * sizeof(cudaStream_t));

    for(int i = 0; i < nStreams; i++)
        cutilSafeCall( cudaStreamCreate(&(streams[i])) );

    
    //병행 실행 시간 측정
    cudaEvent_t StreamStart, StreamStop;
    float StreamTime;

    cudaEventCreate(&StreamStart);
    cudaEventCreate(&StreamStop);

    int offset = 0;

    cudaEventRecord(StreamStart, 0);

    //호스트 디바이스 입력 데이터 전송
    for(int i = 0; i < nStreams; i++)
    {
        offset= i*N/nStreams;
        cudaMemcpyAsync(dev_In + offset, host_In + offset, Size, cudaMemcpyHostToDevice, streams[i]);
    }

    //덧셈 계산
    for(int i = 0; i < nStreams; i++)
    {
        offset= i*N/nStreams;
        kernel<<<nBlocks/nStreams,nThreads, 0, streams[i]>>>(dev_In+offset, dev_Out+offset);
    }

    //디바이스 호스트 출력 데이터 전송
    for(int i = 0; i < nStreams; i++)
    {
        offset= i*N/nStreams;
        cudaMemcpyAsync(host_Out + offset, dev_Out + offset, Size, cudaMemcpyDeviceToHost, streams[i]);
    }

    cudaEventRecord(StreamStop, 0);
    cudaEventSynchronize(StreamStop);

    cudaEventElapsedTime(&StreamTime, StreamStart, StreamStop);


    printf("스트림 실행시간: %f msec\n",StreamTime);
    

    cudaEventDestroy(StreamStart);
    cudaEventDestroy(StreamStop);

    for(int i = 0; i < nStreams; i++)
        cudaStreamDestroy(streams[i]);

    cudaFree(dev_In);
    cudaFree(dev_Out);

    cudaFreeHost(host_In);
    cudaFreeHost(host_Out);

    return 0;
}