불완전한 전역 메모리 액세스가 통합 되었습니까?

Question

n < warpSize 인 경우 병합됩니까? 일부 N이 warpSize로 나눌 경우

// In kernel 
int x; 
if (threadId < n) 
    x = globalMem[threadId]; 

이러한 상황은주기의 마지막 반복에서 appers. 나는이 사이트들에 대해 달리고 장치 메모리를 warpSize으로 나눌 수있을만큼 할당해야합니까, 아니면 그대로 합체 되야합니까?

Answer 1

threadId이 cuda programming guide - thread hierachy에 문서화 된대로 올바르게 계산되는 경우이 액세스가 병합됩니다. threadId = threadIdx.x의 경우입니다.

다른 컴퓨팅 아키텍처의 메모리 병합은 약간 다릅니다. 자세한 내용은 appendix G of cuda programming guide에서 확인할 수 있습니다.

일반적으로 스레드가 첫 번째 스레드가 액세스하는 요소의 주소에서 시작하여 메모리의 연속 요소를 가져 오는 경우 전역 메모리 액세스가 병합된다고 말할 수 있습니다.

플로트 배열이 있다고 가정 해 보겠습니다.
float array[]
및 메모리 액세스를 사용자 액세스 coalesed 것보다 그런 식으로

array[threadIdx.x == 0, threadId.x == 1, threadIdx.x == 2, ..., threadIdx.x == 31] 

에서 찾습니다.

당신이 그런 식으로 메모리에 액세스한다면 (예를 들어 인터리브)

array[threadIdx.x == 0, NONE, threadId.x == 1, NONE, threadIdx.x == 2, ..., NONE, threadIdx.x == 31] 

액세스가 합체

(NONE이 배열 요소가 어떤 스레드가 액세스 할 수 없음을 의미)하지보다 첫 번째 경우에는 연속 128 바이트 메모리를 사용합니다. 두 번째 경우에는 256 바이트를 가져옵니다. 두 번째 경우에는 첫 번째 경우에 하나의 워프 대신 전역 메모리에서 메모리를로드하는 데 두 개의 워프가 필요합니다. 그러나 두 경우 모두 다음 계산을 위해서는 32 개의 float 요소 (128 바이트) 만 필요합니다. 따라서 간단한 경우 글로벌 부하율이 1.0에서 0.5로 떨어집니다.