2.x 디바이스에서 뱅크 충돌

Question

2.x 디바이스가있는 디바이스에서 뱅크 충돌이 무엇입니까? CUDA C 프로그래밍 가이드를 이해함에 따라 2.x 장치에서 두 스레드가 동일한 공유 메모리 뱅크에서 동일한 32 비트 워드를 액세스하면 뱅크 충돌이 발생하지 않습니다. 대신, 그 단어가 방송됩니다. 두 스레드가 동일한 공유 메모리 뱅크에 동일한 32 비트 워드를 쓰면 하나의 스레드 만 성공합니다.

온칩 메모리가 64KB (공유 메모리는 48KB, L1은 16KB 또는 그 반대로)이며 32 개 뱅크로 구성되어 있으므로 각 뱅크는 2KB로 구성됩니다. 따라서 두 스레드가 동일한 공유 메모리 뱅크에서 두 개의 서로 다른 32 비트 워드에 액세스하면 은행 갈등이 발생한다고 생각합니다. 이 올바른지?

Answer 1

설명이 정확합니다. 은행 갈등을 야기 할 수있는 많은 액세스 패턴이 있지만 간단하고 일반적인 예가 있습니다 : strided access.

__shared__ int smem[512]; 

int tid = threadIdx.x; 

x = smem[tid * 2]; // 2-way bank conflicts 
y = smem[tid * 4]; // 4-way bank conflicts 
z = smem[tid * 8]; // 8-way bank conflicts 
// etc. 

은행 ID = 인덱스 % 32는 X, Y의 주소 패턴 보면, Z는 액세스하는 경우, X의 32 개 각각의 스레드 날실이 볼 수 있도록, 2 개 스레드 것 y는 4 개의 스레드가 각 뱅크에 액세스하고, z의 경우 8 개의 스레드가 각 뱅크에 액세스합니다.