커널함수는 비동기 함수이기 때문에, 동기화 처리시, 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;
}

1. 고정메모리.

    - 가상메모리를 사용하지 않고, 메모리에 직접 접근하여 처리하는 접근 방식.

    - 가상 메모리와 물리적 메모리간의 치환 제거.

    - CUDA 스트림을 사용하기 위해서는 고정 메모리는 필수다.

   - 관련 함수 

     1. cudaMallocHost

     2. cudaFreeHost 

2. 제로 메모리.

   - 호스트와 디바이스간에 메모리 복사를 하지 않고, 고정된 메모리[Pinned Memory]영역에 바로 엑세스하여 데이터를 읽고 쓸수 있는 메모리.

   - PCI를 사용해서 전송 속다가 빨라 지는 것은 아니지만.

   - 데이터를 계산하고, 결과값을 메모리에 쓰면 비동기 양방향 PCI 전송이 진행되기 때문에, 그 만큼 성능향상을 볼 수 있다.

   - 전제조건

      - 맵드 메모리를 사용할 때, 글로벌 메모리의 결합 전송과 동일한 조건을 커널에서 충족시켜야 한다. 

       글로벌 메모리 결합 전송조건.

        - 글로벌 메모리를 읽어 올 때 최대 밴드 폭을 사용할 수 있는 조건  - 사실 최신 사양에서는 무시해도 되지 않나 싶다.[개인생각]

    - 주의-

      - 커널에서 작은 크기의 데이터를 많은 횟수로 맵드 메모리를 엑세스 하게 되면, 통상적인 데이터 전송보다 떨어지는 효과를 얻을 수 있다.

   - 관련 함수

      1  cudaHostAlloc

      2  cudaHostGetDevicePointer 

      3 예제: 

3.포터블 고정 메모리.

- 제로 메모리는 하나의 디바이스에서만 유효하기 때문에, 두개의 디바이스상에서 문제가 된다.  즉 두개의 호스트 스레드를 생성하여 처리하게 되는데, 이 때 하나의 스레드에서 생성한 고정된 메모리는 다른 쓰레드에서는 사용할 수 없게 되어, 자원 낭비를 하게 된다.

이런 자원 낭비를 피하고자 사용하는 것이 포터블 고정 메모리라고 한다.  

- 적용.

   cudaHostAlloc() 세번째 파라미터: cudaHostAllocMapped | cudaHostAllocPortable 옵션 지정.

- 메모리 타이별 특성.

  1. 레지스터 

     - 쓰래드 내부에서 몇 개를 사용하는지는 알 수 없다. 다만 툴을 이용해서 확인은 가능.

     - 단항연산을 할수록, 사용하는 레지스터 갯수는 줄어든다.

  2. 로컬메모리

  3. 공유메모리 Shared Memory

       

  4. 글로벌 메모리.

  5. 상수 메모리.

  6. 텍스트 메모리.

-maxrregcount  : 쓰레드당 사용할 수 있는 레지스터 개수를 지정할 수 있다.

                       초과 되었을 경우, 로컬 메모리[DRAM] 를 사용하게 된다. 

1D grid of 1D blocks

__device__ int getGlobalIdx_1D_1D()

{

return blockIdx.x *blockDim.x + threadIdx.x;

}

1D grid of 2D blocks

__device__ int getGlobalIdx_1D_2D()

{

return blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;

}

1D grid of 3D blocks

__device__ int getGlobalIdx_1D_3D()

{

return blockIdx.x * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;

} 

2D grid of 1D blocks 

__device__ int getGlobalIdx_2D_1D()

{

int blockId   = blockIdx.y * gridDim.x + blockIdx.x;

int threadId = blockId * blockDim.x + threadIdx.x; 

return threadId;

}

2D grid of 2D blocks  

 __device__ int getGlobalIdx_2D_2D()

{

int blockId = blockIdx.x + blockIdx.y * gridDim.x; 

int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;

return threadId;

}

2D grid of 3D blocks

__device__ int getGlobalIdx_2D_3D()

{

int blockId = blockIdx.x 

+ blockIdx.y * gridDim.x; 

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)

  + (threadIdx.z * (blockDim.x * blockDim.y))

  + (threadIdx.y * blockDim.x)

  + threadIdx.x;

return threadId;

} 

3D grid of 1D blocks

__device__ int getGlobalIdx_3D_1D()

{

int blockId = blockIdx.x 

+ blockIdx.y * gridDim.x 

+ gridDim.x * gridDim.y * blockIdx.z; 

int threadId = blockId * blockDim.x + threadIdx.x;

return threadId;

} 

3D grid of 2D blocks

__device__ int getGlobalIdx_3D_2D()

{

int blockId = blockIdx.x 

        + blockIdx.y * gridDim.x 

+ gridDim.x * gridDim.y * blockIdx.z; 

int threadId = blockId * (blockDim.x * blockDim.y)

 + (threadIdx.y * blockDim.x)

 + threadIdx.x;

return threadId;

}

3D grid of 3D blocks

__device__ int getGlobalIdx_3D_3D()

{

int blockId = blockIdx.x 

+ blockIdx.y * gridDim.x 

+ gridDim.x * gridDim.y * blockIdx.z; 

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)

 + (threadIdx.z * (blockDim.x * blockDim.y))

 + (threadIdx.y * blockDim.x)

 + threadIdx.x;

return threadId;

}

1. visual Studio NSIGHT 메뉴에 Start Performance Analysis 항목을 선택. 

2. Activity Type 에서 Profile CUDA Application 항목을 선택한다.

3. Application Control 항목에서 Launch 버튼을 클릭한다.

4. Launch 버튼을 클릭하면, kill 버튼을 클릭할 때까지 반복해서 프로그램을 실행한다.

5. Kill 을 클릭하면, 리포팅을 확인 할 수 있다.

CUDA  디버깅 환경 구성.

- 프로그램 설치. [Nsight Visual Studio Edition 5.2 버전 기준]

- Nsight Monitor 

  4. Disable D3D acceleration for WPF

    - DisableWpfHardwareAcceleration.reg 항목을 검색하여, 등록한다. [C:\Program Files (x86)\NVIDIA Corporation\Nsight Visual Studio Edition 5.2\Host\Common ] 

 5. 컴퓨터 재부팅.

- Visual Studio 2013 

1. Nsight User Properties 설정

 1. 솔류션 검색창->프로젝트 선택 후 팝업메뉴에서 Nsight User Properties 항목을 선택한다.      

2. Launch 항목을 선택한다.

2. Run the Memory Checker 체크.

:: CUDA Visual Profiler 

  - 소 개

  - 실 행

    1. console

bank conflict

SIMT를 실행시킬때 문제중의 하나가 memory access이다. GPU에서는 동시에 여러개의 데이터를 처리해야하기 때문에, 동시에 여러개의 데이터에 access를 허용한다. 이것을 하기 위해서 GPU는 shared memory를 각 warp마다 일정 갯수의 memory bank로 나누어 두었는데, 각각의 bank는 bank단위로 동시에 접근할 수 있다. 이때 bank conflict란 프로그래밍 잘못으로 동시에 서로 다른 thread가 특정 bank를 access할때 발생하는 문제이다.

서로 다른 thread가 하나의 특정 bank에 access하게 되면, 각각의 thread는 해당 bank에 접근하기 위해서 순차적으로 변하게 되고, 이는 병렬적으로 처리하려고 했던 의도를 벗어나게된다. 즉, 의도하지 않은 행동이 된다는 것이다.

이것은 GPU를 이용한 병렬처리에서 속도를 낮추게 하는 문제를 일으킨다. 여기서는 해당 문제 상황에 대해서 정리를 하고자 한다.

1.x대에서는 warp의 크기가 32이고, bank의 갯수가 16개이다. 이것은 최대 shared memory에 접근할 수 있는 thread의 갯수는 16개뿐이라는 이야기다. 이 문제를 해결하기 위해서 CUDA는 warp를 두개의 half-warp로 나눈다. 두개의 half-warp를 순차적으로 실행하므로써 bank conflict가 사라진다.(단, shared memory access를 하는데, 2 cycle이 걸리겠지?)

bank conflict가 발생하는 다른 경우를 살펴보자.

1. 다음과 같은 코드가 있다고 하자.

__shared__ struct type shared[32]; 

struct type data = shared[BaseIndex + tid];

이때

1) type 구조체가 다음과 같이 정의되었다면,

struct type { 
float x, y, z; 
};

위의 구조체에 의해서 실행되는 thread들은 bank conflict를 발생시키지 않는다. 
여기서 bank conflict가 발생하지 않는 이유를 설명하기 위해서, bank가 2개이고, thread도 2개라고 가정을 하자. 또한 데이터는 순서대로 bank 0과 bank 1을 왔다갔다하면서 쌓인다고 하자. 그럼 다음과 같이 데이터들이 저장되어 있을 것이다. (bank의 access단위는 32bits라고 하자.)

이때 위의 구조체대로라면, Thread 0은 [0]을 Thread 1은 [3]에 접근한다. [0]은 Bank 0에 있고, [3]은 Bank 1에 있으므로, 다른 Thread가 동시에 한 Bank에 접근하지 않는다. 

그러므로 Bank Conflict가 발생되지 않는다.

2) 만약 구조체가 다음과 같다면,

struct type { 
float x, y; 
};

위의 코드는 bank conflict를 발생시킨다. 이유는 위의 1)번과 같은 가정을 했을때, thread 0은 [0], thread 1은 [2]에 접근하는데, 이때 [0]과 [2]는 table에서 보듯이 같은 bank에 있으므로, 두개의 thread가 동시에 같은 bank를 access하게 된다. 그러므로 bank conflict가 발생한다.

예제 1

__shared__ char shared[32]; 
char data = shared[BaseIndex + tid];

1. 결과: 위의 코드는 Bank conflict를 일으킨다. 

2. 원인: thread 0은 shared[0]을 (BaseIndex는 0이라 하자.), thread 1은 shared[1]에 접근하는데, 이때 char는 1byte이고, bank의 한 데이터는 4bytes이므로, bank0에서 [0]및 [1]을 읽어야 한다. 즉, 같은 bank를 access하므로 bank conflict가 발생한다.

3.해결:

__shared__ char shared[32]; 
char data = shared[BaseIndex + 4 * tid];

와 같이 하면된다.

예제 2

__shared__ double shared[32]; 
double data = shared[BaseIndex + tid];

1. 결과: 위의 코드는 Bank conflict를 일으킨다.

2. 원인: thread 0은 shared[0]을 (BaseIndex는 0이라 하자.), thread 1은 shared[2]에 접근한다. 이유는 double은 8bytes이므로 한 bank에서만 읽어서는 안된다. 이와 같은 이유로 같은 bank를 access하므로 bank conflict가 발생한다.

3. 결과: 쪼개서 읽어야겠다.  어떻게 쪼개지?

예제 3 s는 stride이다.

__shared__ float shared[32]; 
float data = shared[BaseIndex + s * tid];

위와 같은 코드가 있을 때 각 thread가 s가 어떤 조건을 가져야지만, bank conflict가 나지 않을지를 찾아보자. 
bank갯수가 m이라고 하고, 위의 요청 갯수가 n개라고 하자.

이때 bank conflict가 일어나지 않을 조건은 s*n이 m의 배수이거나 m과 같을 경우(예, 위의 1)에 조건을 기준으로 설명을 하면, ), m과 s의 최대공약수가 d라고 하면, n이 m/d의 배수일 경우이다.

출처: https://dyanos.wordpress.com/2009/10/21/cuda-%ED%94%84%EB%A1%9C%EA%B7%B8%EB%9E%98%EB%B0%8D-%EC%A0%95%EB%A6%AC/  에서 정리 함. [임시보관]

2차원 배열일 때.