CUDA에서 공유 메모리 뱅크 충돌 : 메모리가 뱅크에 정렬되는 방식

Question

공유 메모리는 뱅크로 나누어지고 동일한 스레드 내의 단일 데이터 요소에 대한 다중 스레드에 의한 액세스는 충돌 (또는 브로드 캐스트). N 쌍의 수와 한 쌍의 제 16 개 플로트가

__shared__ float A[34*N] 

한 매트릭스이며, 다음 순간

는 I 개념적 두 행렬의 여러 쌍을 나타내는 상당히 큰 배열을 할당 18 개의 수레가 두 번째입니다.

사실, 첫 번째 행렬에 대한 액세스는 충돌이 없지만 두 번째 행에 대한 액세스에는 충돌이 있습니다. 이러한 갈등은 피할 수없는 일이지만, 제 생각에는 두 번째 행렬이 18이므로 미래의 모든 행렬이 은행과 정렬되지 않아 필요한 것보다 더 많은 갈등이 발생할 것입니다.

그렇다면 어떻게해야합니까?

공유 메모리를 할당 할 때마다 새 뱅크에서 시작합니까? 그래서 내가 할 수있는 잠재적으로 할 수

__shared__ Apair1[34] 
__shared__ Apair2[34] 
... 

어떤 아이디어?

감사 행렬의 당신의 쌍 연속적으로 저장되고, 당신이 스레드 인덱스로 선형 적 요소를 액세스하는 경우, 당신은 메모리 뱅크 충돌을 공유하지 않을 경우

Answer 1

.

A[0] <- mat1 element1 
A[1] <- mat1 element2 
A[2] <- mat1 element3 
A[15] <- mat1 element16 
A[16] <- mat2 element1 
A[17] <- mat2 element2 
A[33] <- mat2 element18 

그리고 당신이 사용하여 액세스 : 즉

당신이있는 경우

float element; 
element = A[pairindex * 34 + matindex * 16 + threadIdx.x]; 

그리고 인접한 스레드가 매트릭스에 인접한 요소에 액세스하면 충돌이 없습니다.

귀하의 의견 (아래)에 대한 응답으로 귀하가 귀하의 이해를 잘못 알고있는 것처럼 보입니다. 16 개의 뱅크 (현재 세대에서는 32 개, 차세대에서는 32 개, 페르미 (Fermi))가 있지만, 연속하는 32 비트 워드가 연속 뱅크에 상주한다는 점은 사실이다. 즉, 어드레스 공간은 뱅크간에 인터리빙된다. 즉, 배열 인덱스가 항상 x + threadIdx.x (여기서 x은 threadIdx.x에 종속적이지 않거나 최소한 16 개의 스레드 그룹간에 일정 함)로 분해 될 수 있다면 배열 충돌을 방지 할 수 있습니다.

배열을 따라 더 멀리 매트릭스에 액세스하면 여전히 인접한 청크에서 액세스하므로 뱅크 충돌이 발생하지 않습니다. 은행 갈등이있는 인접하지 않은 요소에 액세스하기 시작할 때만입니다.

은 SDK의 감소 샘플, 최적화 된 구현 순진 구현에서 살펴 본다 가능 가치 구축하여 잘 은행 충돌을 보여줍니다.

Answer 2

연속적인 32 비트가 다음 뱅크에 있도록 뱅크가 설정됩니다. 따라서 4 바이트 부동 소수점 배열을 선언하면 배열의 각 후속 부동 소수점은 다음 뱅크에있게됩니다 (아키텍처에 따라 모듈로 16 또는 32). 저는 여러분이 컴퓨팅 능력 1.x에 있다고 가정 할 것이므로 너비가 16 인 뱅크를가집니다.

배열이 18과 16 인 경우 재미있을 수 있습니다. 당신은 (당신이 충돌을 받고 있다면 나는 당신이하고있는 가정으로) threadIdx.x를 사용하여 트랜스 요소에 액세스 할 때 은행 충돌을 피할 수

__shared__ float sixteen[16][16+1] 

처럼 선언함으로써 16 × 16 배열의 은행 충돌을 피할 수 있습니다. 예를 들어 16x16 매트릭스의 첫 번째 행에있는 요소에 액세스 할 때, 모든 요소는 첫 번째 뱅크에 상주합니다. 당신이하고 싶은 일은 이들 각각을 연속적인 은행에 두는 것입니다. 패딩이 당신을 위해 이것을합니다. 16 행 [열] [열], 또는 평평한 행렬에 대해 마찬가지로 정확히 16 행 (16 + 1) + 열로 배열을 처리합니다.

18x18의 경우, 조바꿈에 액세스 할 때 평평한 움직임을 보입니다. 그 대답은 다시는 (첫 번째 열의 요소를 액세스하는 말) 전치에 액세스 할 때 (18 + 1) % 16 = 3으로 접속되며, 이제

__shared__ float eighteens[18][18+1] 

(1)에 의해 패드이고, 3, 6, 9, 12, 15, 2, 5, 8 등의 은행에 액세스하므로 아무런 충돌이 없어야합니다.

크기가 18 인 행렬을 사용하기 때문에 발생하는 특정 정렬 이동은 문제가되지 않습니다. 왜냐하면 배열의 시작점은 아무런 차이가 없으므로 사용자가 액세스하는 순서 일뿐입니다. 위에서 제안한 배열을 병합하고 1로 병합하려면 비슷한 방식으로 액세스하는 한 괜찮습니다.