인프런 커뮤니티 질문&답변
작성한 질문수
global memory를 사용한 matrix copy 대비 global memory를 활용한 matrix transpose가 속도가 느린 이유가 궁금합니다.
해결된 질문
작성
·
243
0
안녕하세요.
먼저 좋은 강의 감사드립니다.
아래 소스는 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 때문으로 추정하고 있는데, 맞는 생각일까요?
감사합니다.
답변 1
0
안녕하세요.
이후 강의에서도 계속 나오게 됩니다만, 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++ 코딩에도 해당되는 내용입니다.
감사합니다.
상세하고 빠른 답변 감사드립니다.