임시 - 쿠다 공유메모리 뱅크 충돌이란.

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라고 하자.)

Bank 0	Bank 1
[0]	[1]
[2]	[3]
[4]	[5]
…	…

이때 위의 구조체대로라면, 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차원 배열일 때.

//뱅크갯수가 32개 가정.

__global__ void BankConflict(float* gData)

{

    __shared___ float sData[32[32];       

    float Data = sData[threadIdx.x][0];

}   

0 번쓰레드 -> 0번뱅크 0번째.

1 번쓰레드 -> 0번뱅크 1번째.

2 번쓰레드 -> 0번뱅크 2번째.

3 번쓰레드 -> 0번뱅크 3번째.

4 번쓰레드 -> 0번뱅크 4번째.

1. 결국 뱅크 충돌.

해결. 두번째 열를 1을 더해 33개로 하여, 한나씩 밀리도록 한다.

0 번쓰레드 -> 0번뱅크 0번째.

1 번쓰레드 -> 0번뱅크 32번째.

2 번쓰레드 -> 1번뱅크 31번째.

3 번쓰레드 -> 2번뱅크 30번째.

4 번쓰레드 -> 3번뱅크 29번째.

이와 같이 1열을 더 추가하여- sData[32][33] -, 뱅크충돌을 회피하는 방법을 메모리 패딩이라고 한다.