global memory를 사용한 matrix copy 대비 global memory를 활용한 matrix transpose가 속도가 느린 이유가 궁금합니다.

Question

안녕하세요. 먼저 좋은 강의 감사드립니다. 아래 소스는 matcpy-dev.cu 에서 가져온 kernel 함수입니다. global void kernelMatCpy( float* C, const float* A, int matsize, size_t pitch_in_elem ) { register unsigned gy = blockIdx.y * blockDim.y + threadIdx.y; // CUDA-provided index if (gy < matsize) { register unsigned gx = blockIdx.x * blockDim.x + threadIdx.x; // CUDA-provided index if (gx < matsize) { register unsigned idx = gy * pitch_in_elem + gx; // in element C[idx] = A[idx]; } } } 그리고 다음 소스는 transpose-dev.cu 에서 가져온 kernel 함수 입니다. global void kernelMatTranspose( float* C, const float* A, unsigned matsize, size_t pitch_in_elem ) { register unsigned gy = blockIdx.y * blockDim.y + threadIdx.y; // CUDA-provided index if (gy < matsize) { register unsigned gx = blockIdx.x * blockDim.x + threadIdx.x; // CUDA-provided index if (gx < matsize) { register unsigned idxA = gy * pitch_in_elem + gx; register unsigned idxC = gx * pitch_in_elem + gy; C[idxC] = A[idxA]; } } } 메모리 접근 관점에서 보면 각 thread별로 A에서 read해서 C에 write하는 동일한 과정으로 보이는데, 결과는 속도 차이가 많이 나는 정확한 원인이 궁금합니다. memory coalescing 때문으로 추정하고 있는데, 맞는 생각일까요? 감사합니다.

onemoresipofcoffee · Answer

안녕하세요. 이후 강의에서도 계속 나오게 됩니다만, CUDA는 물론이고, 일반적인 C, C++ 코딩에서도 memory coalescing 문제는 전체 시스템의 성능을 대폭 떨어뜨립니다. 위 예제에서도, transpose 연산이 얼핏 보기에는 matrix 원소 하나 하나를 재배치하는 것 뿐이라고 보이지만, 실제로는 main memory 에서든, CUDA 의 global memory 에서든, memory access 시에는 memory access 순서에 따라서, 성능 차이가 꽤 나게되고, 예를 들어, A번지, A+1번지, A+2 번지 처럼, 연속된 메모리 영역을 access 하는 것과, 꽤 떨어진 메모리 영역을 access 하는 것은 cache fail 부터, DRAM 의 구조 등으로 성능 차이가 꽤 납니다. CUDA 코딩의 많은 부분은 메모리 access 패턴을 최적화하는 부분을 포함하게 되고, 이후 동영상에서도 많은 설명이 나오게 됩니다. 이부분은 CUDA 만이 아니고, 기존의 C, C++ 코딩에도 해당되는 내용입니다. 감사합니다.