Shared Memory (2) - Square/Rectangular Shared Memory

References

• Professional CUDA C Programming

Contents

• Checking the Data Layout of Shared Memory
• Square Shared Memory
• Rectangular Shared Memory

---

지난 포스팅에 이어서 공유 메모리에 대한 내용에 대해서 더 알아보도록 하겠습니다.

Share Memory (1)

Share Memory (1)

 

Checking the Data Layout of Shared Memory

 

shared memory를 효과적으로 사용하는 방법에 대해서 알아보기 위해, 이번 포스팅과 다음 포스팅에서는 다음의 몇 가지 주제들을 간단한 예제를 통해 설명하도록 하겠습니다.

• Square versus rectangular arrays
• Row-major versus column-major accesses
• Static versus dynamic shared memory declarations
• File-scope versus kernel-scope shared memory
• Memory padding versus no memory padding

 

shared memory를 사용하는 커널을 디자인할 때에는 아래의 두 가지 컨셉에 주목해야 합니다.

• Mapping data elements across memory banks
• Mapping from thread index to shared memory offset

위 두 가지 개념이 명확하다면 bank conflict를 피하고 shared memory의 이점을 완전히 활용할 수 있는 커널을 디자인할 수 있습니다.

---

Square Shared Memory

shared memory를 사용하여 정사각형 차원의 global data를 캐싱할 수 있습니다. 정사각형 배열의 심플한 차원은 2D 스레드 인덱스로부터 1D 메모리 오프셋을 쉽게 계산할 수 있습니다. 아래 그림은 각 차원이 32개의 원소인 shared memory tile을 보여줍니다.

위 이미지 상단은 1D 데이터 레이아웃의 실제 정렬을 보여주고, 하단은 논리적 2차원 shared memory view를 보여줍니다.

 

2차원 shared memory 변수는 다음과 같이 동적으로 선언할 수 있습니다.

__shared__ int tile[N][N];

이 shared memory tile은 정사각형이기 때문에, x 또는 y 차원의 인접한 요소에 액세스하는 이웃 스레드들의 2D 스레드 블록으로부터 액세스할 수 있습니다.

tile[threadIdx.y][threadIdx.x]
tile[threadIdx.x][threadIdx.y]

위의 두 가지 방법 중에 어떤 것이 더 좋을까요?

이를 알기 위해서는 어떻게 스레드가 shared memory banks에 매핑하는지에 주목해야합니다. 이전 포스팅에서 동일한 warp 내의 스레드들은 분리된 banks에 액세스하여 최적화한다고 하였습니다. 동일한 warp의 스레드들은 threadIdx.x의 연속적이 값으로 식별됩니다. 다른 bank에 속하는 shared memory의 원소들 또한 연속적으로 저장됩니다. 그러므로, threadIdx.x의 연속적인 값을 갖는 스레드들이 shared memory의 연속적인 위치에 액세스하는 것이 가장 좋습니다.

따라서, 인접한 스레드가 인접한 배열 셀에 액세스하기 때문에 tile[threadIdx.y][threadIdx.x] 패턴이 좋은 성능을 보여주고 두 번째 패턴보다 bank conflict가 덜 발생한다는 것을 알 수 있습니다. 

 

Accessing Row-Major versus Column-Major

각 차원에서 32개의 스레드를 포함하는 하나의 2D 블록의 그리드를 사용하는 예제를 살펴보겠습니다. 블록의 차원은 다음의 macro로 정의합니다.

#define BDIMX 32
#define BDIMY 32

이 매크로를 사용하여 커널의 execution configuration은 다음과 같이 정의할 수 있습니다.

dim3 block(BDIMX, BDIMY);
dim3 grid(1,1);

커널은 다음의 두 가지 간단한 동작을 수행합니다.

• global thread 인덱스를 row-major 순서로 2D shared memory array에 write
• row-major 순서로 shared memory로부터 값들을 읽고, global memory에 저장

따라서, 이 커널은 다음과 같이 작성할 수 있습니다.

__global__
void setRowReadRow(int* out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.y][threadIdx.x];
}

이 커널에서는 다음의 3가지 메모리 동작이 수행됩니다.

• One store operation on shared memory (line 11)
• One load operation on shared memory (line 17)
• One store operation on global memory (line 17)

동일한 warp의 스레드들은 연속적인 threadIdx.x 값을 가지고 있고, shared memory 배열인 tile의 가장 내부 차원을 인덱싱하기 위해 threadIdx.x를 사용하기 때문에 이 커널은 bank conflict가 발생하지 않습니다.

 

반면에, shared memory tile에 data 데이터를 할당할 때 threadIdx.y와 threadIdx.x를 서로 바꾸면, 한 워프에서의 메모리 액세스는 column-major 순이 됩니다. 모든 shared memory load와 store에서 bank conflict를 발생합니다.

__global__
void setColReadCol(int* out)
{
    // static shared memory
    __shared__ int tile[BDIMX][BDIMY];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.x][threadIdx.y] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

 

커널을 실행하기 위한 메인 함수는 아래 링크를 참조해주세요.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/SharedMemory/smemSquare.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

 

위 파일을 컴파일하고 nvprof로 프로파일링 해보도록 하겠습니다.

제가 사용하는 GPU는 4-byte shared memory access mode가 default로 사용되고 있습니다.

커널의 수행 시간을 살펴보면 row-major가 column-major보다 약 2배 가량 빠릅니다. 이는 인접한 스레드가 인접한 words에 액세스하기 때문입니다.

 

이제 두 커널 함수의 bank conflict를 체크해보도록 하겠습니다. nvprof를 사용한다면 아래의 메트릭을 사용하면 측정할 수 있습니다만, 

compute capability 7.5 이상에서는 메트릭을 사용할 수 없기 때문에 nsight compute으로 측정해보도록 하겠습니다.

하지만, 위와 관련된 '*_per_request'는 제공되지 않아서, shared_load_transactions와 shared_store_transactions를 대신 측정하였습니다.

ncu.bat --metrics l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum,l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum ./smemSquare.exe

결과는 위와 같습니다. column-major인 setColReadCol의 load/store transcation 요청은 총 1024번이지만, row-major인 setRowReadRow는 32번밖에 되지 않습니다.

 

Writing Row-Major and Reading Column-Major

이번에는 shared memory write는 row-major로 수행하고, shared memory read는 column-major로 수행하는 커널을 작성해보겠습니다.

__global__
void setRowReadCol(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

위 커널 함수를 추가하고 컴파일 후 다시 shared memory의 load/store transcation를 측정해보도록 하겠습니다.

store operation은 bank conflict가 발생하지 않았지만, load operation은 bank conflict가 발생했습니다.

 

Dynamic Shared Memory

다음으로 동적으로 shared memory를 선언한 동일한 커널들을 실행해보도록 하겠습니다.

동적 shared memory는 커널 외부에 선언하여 file scope에서 global로 만들거나 커널 내에 선언하여 kernel scope를 가지도록 할 수 있습니다. Dynamic shared memory는 반드시 크기가 지정되지 않은 1차원 배열로 선언되어야 합니다. 그러므로 2차원 스레드 인덱스를 기반으로 메모리 액세스 인덱스를 계산해야합니다.

커널 내에서 row-major 순서로 write하고 column-major 순서로 read하도록 작성하기 때문에 아래의 2개의 인덱스를 사용하도록 하겠습니다.

 

Dynamic Shared Memory를 사용한 커널은 다음과 같습니다.

__global__
void setRowReadColDyn(int *out)
{
    // dynamic shared memory
    extern  __shared__ int tile[];

    // mapping from thread index to global memory index
    unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;

    // shared memory store operation
    tile[row_idx] = row_idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[row_idx] = tile[col_idx];
}

그리고, shared memory 크기는 커널을 실행할 때 아래와 같이 지정해주어야 합니다.

setRowReadColDyn<<<grid, block, BDIMX*BDIMY*sizeof(int)>>>(d_C);

 

위 커널을 추가하고, 다시 프로파일링해보도록 하겠습니다.

결과는 이전의 setRowReadCol 커널과 완전히 동일합니다. 

 

Padding Statically Declared Shared Memory

이전 포스팅에서 Memory Padding에 대해서 이야기를 했었습니다. 이렇게 bank conflict가 발생할 때, memory padding을 추가해주면 conflict를 피할 수 있다고 언급했습니다. 간단하게 다음과 같이 2차원 shared memory에 1개의 column을 추가하여 setRowReadCol 커널과 동일한 커널을 작성하고 프로파일링 해보도록 하겠습니다.

__global__
void setRowReadColPad(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX + IPAD];

    // mapping from thread index to global memory offset
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

위 커널을 추가하고 프로파일링한 결과는 다음과 같습니다.

padding을 하나 추가함으로써 bank conflict 문제를 해결하였습니다.

(다만, Kepler는 64-bit access mode를 사용하므로 이것으로 항상 해결할 수 있는 것은 아닙니다.)

 

Padding Dynamically Declared Shared Memory

dynamic shared memory에 padding하는 것은 조금 더 복잡합니다. 2차원 인덱스로부터 1차원 메모리 인덱스를 변환할 때, 각 row에서 하나의 padding된 메모리를 스킵해주어야 합니다.

다음의 그림은 간단하게 5개의 bank로 이루어진 shared memory를 사용하여 인덱스를 계산하는 것을 보여줍니다.

Dynamic Shared Memory에 padding을 추가하는 커널은 다음과 같습니다.

__global__
void setRowReadColDynPad(int *out)
{
    // dynamic shared memory
    extern  __shared__ int tile[];

    // mapping from thread index to global memory index
    unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;
    unsigned int col_idx = threadIdx.x * (blockDim.x + IPAD) + threadIdx.y;

    unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[row_idx] = g_idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[g_idx] = tile[col_idx];
}

data를 저장하는데 사용되는 global memory는 padded shared memory보다 작기 때문에, 3개의 인덱스가 필요합니다. 하나는 shared memory에 write를 하기 위한 row-major 인덱스이고, 하나는 shared memory로부터 read하기 위한 column-major 인덱스, 나머지 하나는 unpadded global memory에 coalesced accesses를 위한 인덱스입니다.

 

위 커널은 다음과 같이 실행될 수 있습니다.

setRowReadColDynPad<<<grid, block, (BDIMX + IPAD)*BDIMY*sizeof(int)>>>(d_C);

다시 nsight compute로 측정하면 다음과 같은 결과를 확인하실 수 있습니다.

동적으로 선언된 shared memory에 padding한 것과 동일한 결과를 보여줍니다.

 

Comparing the Performance of the Square Shared Memory Kernels

위에서 살펴본 모든 커널들의 실행 시간을 살펴보겠습니다.

결과로부터

• padding을 사용한 커널은 감소된 bank conflict 때문에 좋은 성능을 보여줍니다.
• 동적으로 shared memory를 선언한 커널은 overhead가 적습니다.

라는 것을 확인할 수 있습니다.

 

---

Rectangular Shared Memory

이번에는 위에서 구현하고 비교했던 것들을 Rectangular Shared Memory에서 동일하게 진행해보겠습니다.

직사각형의 shared memory는 더 일반적인 2D shared memory 케이스이며, 배열에서 행과 열의 수가 같지 않습니다. 이제, 위에서 본 정사각형처럼 간단하게 thread 좌표를 변환할 수 없습니다. 위와 같은 방법으로 한다면 rectangular shared memory를 사용할 때 메모리 액세스 violation이 발생할 수 있습니다. 따라서, 위에서 구현한 커널들에서 액세스 인덱스를 새로 계산하도록 하겠습니다.

 

위에서 살펴본 것과 비교하기 위해서 rectangular shared memory 배열을 row당 32개의 요소, column당 16개의 요소로 설정하도록 하겠습니다. 이는 이전과 동일하게 매크로로 설정됩니다.

#define BDIMX 32
#define BDIMY 16

 

Accessing Row-Major versus Accessing Column-Major

정사각형 shared memory의 처음 두 커널과 동일한 커널을 우선 사용합니다.

__global__
void setRowReadRow(int* out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.y][threadIdx.x];
}

__global__
void setColReadCol(int* out)
{
    // static shared memory
    __shared__ int tile[BDIMX][BDIMY];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.x][threadIdx.y] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

 

그리고 nsight compute로 shared_load_transactions과 shared_store_transactions를 측정해보도록 하겠습니다.

결과는 다음과 같습니다.

setRowReadRow 커널은 shared memory에 대한 load/store 요청이 한 번의 transaction으로 처리됩니다. 반면에 setColReadCol 커널은 동일한 요청을 8번의 transactions로 처리합니다.

 

Writing Row-Major and Reading Column-Major

이번에는 row-major로 shared memory write를 수행하고 column-major로 shared memory로부터 read하는 커널을 살펴보겠습니다. 이 커널은 read-world 어플리케이션에서 행렬 transpose를 하는데 사용될 수 있습니다.

 

2D shared memory tile은 다음과 같이 선언됩니다.

__shared__ int tile[BDIMY][BDIMX];

위에서 살펴본 것과 마찬가지로 커널에는 3가지 메모리 operation이 있습니다.

• Write to a shared memory row with each warp to avoid bank conflicts.
• Read from a shared memory column with each warp to perform a matrix transpose.
• Write to a global memory row from each warp with coalesced access.

올바르게 shared memory와 global memory에 액세스하기 위한 과정은 다음과 같습니다.

먼저, 현재 스레드의 2D 스레드 인덱스를 1차원 global thread ID로 변환합니다.

unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

이 idx는 1차원의 row-major 매핑이며, 이는 global memory accesses가 병합되도록 해줍니다. output global memory에서 데이터 원소들이 transposed 되므로, transpose matrix에서의 새로운 좌표를 계산해야 합니다.

unsigned int irow = idx / blockDim.y;
unsigned int icol = idx % blockDim.y;

그리고 global thread ID를 2차원 shared memory tile에 저장하여 shared memory tile을 초기화합니다.

tile[threadIdx.y][threadIdx.x] = idx;

여기서, shared memory의 데이터는 0부터 BDIMX x BDIMY - 1까지의 수로 저장됩니다. 각 warp는 shared memory에 대해 row-major write를 수행하기 때문에 write operation에는 bank conflict가 발생하지 않습니다.

 

이제 irow와 icol을 사용하여 shared memory에 액세스하여 전치된 데이터를 global memory에 1차원 thread ID를 사용하여 write 합니다. 아래 코드에서 warp는 shared memory의 한 column으로부터 요소들을 읽고, global memory에 병합된 write를 수행합니다.

out[idx] = tile[icol][irow];

 

전체 커널은 다음과 같습니다.

__global__
void setRowReadCol(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from 2D thread index to linear memory
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // convert idx to transposed coordinate (row, col)
    unsigned int irow = idx / blockDim.y;
    unsigned int icol = idx % blockDim.y;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[icol][irow];
}

위 커널의 프로파일링 결과는 다음과 같습니다.

store operation은 conflict가 발생하지 않았지만, load operation은 conflict가 발생하였습니다.

 

Dynamically Declared Shared Memory

dynamic shared memory는 오직 1차원 배열로만 선언할 수 있습니다. 따라서 rows로 write를 수행하고, column으로 read를 수행할 때, 2D thread 좌표를 1D shared memory 인덱스로 변환하기 위해서 아래의 새로운 인덱스(icol, irow)가 필요합니다.

unsigned int col_idx = icol * blockDim.x + irow;

icol은 스레드 블록의 가장 안쪽의 차원에 대응되고, 이 변환으로 shared memory에 대한 column-major 액세스가 이루어지고 bank conflict를 발생시킵니다.

커널 코드는 다음과 같습니다.

__global__
void setRowReadColDyn(int *out)
{
    // dynamic shared memory
    extern  __shared__ int tile[];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // convert idx to transposed (row, col)
    unsigned int irow = idx / blockDim.y;
    unsigned int icol = idx % blockDim.y;

    // convert back to smem idx to access the transposed element
    unsigned int col_idx = icol * blockDim.x + irow;

    // shared memory store operation
    tile[idx] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[col_idx];
}

Share Memory가 동적으로 할당되기 때문에 위 커널을 실행할 때에는 다음과 같이 공유 메모리 크기를 execution configuration에 지정해주어야 합니다.

setRowReadColDyn<<<grid, block, BDIMX*BDIMY*sizeof(int)>>>(d_C);

커널을 실행하고, shared memory의 load/store transaction의 총합을 측정해보겠습니다.

write operation에서는 conflict가 발생하지 않았지만, read operation에서는 conflict가 발생한 것을 확인할 수 있습니다. 이처럼 동적으로 shared memory를 할당하는 것은 bank conflict에 영향을 미치지 않습니다.

 

Padding Statically Declared Shared Memory

Rectangular shared memory에서 padding을 사용하면 bank conflict를 해결할 수 있습니다. Square Shared Memory처럼 padding을 추가해주는 이번에는 padding 크기를 2로 설정해보겠습니다.

#define IPAD 2

정적으로 padding을 추가하면, 다음과 같이 shared memory가 선언되어야 합니다.

__shared__ int tile[BDIMY][BDIMX + IPAD];

이제 padding을 사용하는 것을 제외하고 나머지 코드는 setRowReadCol 커널 함수와 동일한 setRowReadColPad 커널 함수를 구현합니다.

__global__
void setRowReadColPad(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX + IPAD];

    // mapping from 2D thread index to linear memory
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // convert idx to transposed (row, col)
    unsigned int irow = idx / blockDim.y;
    unsigned int icol = idx % blockDim.y;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[icol][irow];
}

마찬가지로 위 커널을 추가하여 컴파일 후, shared memory의 load/store transaction을 측정해보겠습니다.

conflict가 하나도 발생하지 않을 것을 확인할 수 있습니다.

 

Padding Dynamically Declared Shared Memory

padding 기법은 dynamic shared memory 커널에도 적용할 수 있습니다. padded shared memory와 global memory는 서로 다른 크기이기 때문에, 스레드당 3개의 인덱스가 커널 내에서 사용되어야 합니다.

• row_idx: padded shared memory의 row index
• col_idx: padded shared memory의 column index
• g_idx: linear global memory의 index

위의 값들은 다음와 같이 계산될 수 있습니다.

// mapping from thread index to global memory index
unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;

// convert idx to transposed (row, col)
unsigned int irow = g_idx / blockDim.y;
unsigned int icol = g_idx % blockDim.y;

unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;

// convert back to smem idx to access the transposed element
unsigned int col_idx = icol * (blockDim.x + IPAD) + irow;

 

전체 커널 함수 setRowReadColDynPad는 다음과 같습니다.

__global__
void setRowReadColDynPad(int *out)
{
    // dynamic shared memory
    extern  __shared__ int tile[];

    // mapping from thread index to global memory index
    unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;

    // convert idx to transposed (row, col)
    unsigned int irow = g_idx / blockDim.y;
    unsigned int icol = g_idx % blockDim.y;

    unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;

    // convert back to smem idx to access the transposed element
    unsigned int col_idx = icol * (blockDim.x + IPAD) + irow;

    // shared memory store operation
    tile[row_idx] = g_idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[g_idx] = tile[col_idx];
}

 

동일하게 shared memory의 load/store transcation의 측정 결과는 다음과 같으며, setRowReadColPad와 동일합니다.

 

Rectangular Shared Memory Kernels의 성능

위에서 살펴본 Rectangular Shared Memory를 사용한 커널들의 수행 시간을 살펴보겠습니다.

row-major write와 column-major read를 수행하는 커널들 중에서는 일반적으로 shared memory padding을 사용한 커널이 bank conflict를 제거하여 더 좋은 성능을 얻는 것을 볼 수 있습니다.

 

---

위에서 사용된 코드는 아래의 링크에서 확인하실 수 있습니다 !

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/SharedMemory/smemSquare.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/SharedMemory/smemRectangular.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread