Tiled 2D Convolution

References

• Programming Massively Parallel Processors

Contents

• Tiled 2D Convolution with Halo Cells

---

1D Convolution (CUDA Constant Memory)

이전 포스팅 1D 컨볼루션에 이어서 이번 포스팅에서는 2D 컨볼루션에 대해 알아보겠습니다. 기본적인 2D 컨볼루션에 관한 것은 이전 포스팅 참조 부탁드립니다. 아래 포스팅에서의 이미지 Blur 처리 예제도 도움이 될 듯 합니다.. !

CUDA Thread 구조와 Data Mapping (예제 : 이미지 흑백, Blur 처리)

 

실제 이미지를 표현할 때 보통 2D-행렬로 표현됩니다. 이미지 처리 라이브러리는 일반적으로 이미지를 메모리로 읽을 때, row-major 형태로 저장합니다. 만약 바이트 단위로 이미지의 너비(width)가 DRAM 버스트(burst) 크기의 배수가 아닌 경우, 1행(row 1)의 시작점이 DRAM 버스트 경계에서 잘못 정렬될 수 있습니다. 이러한 misalignment(정렬 불량)으로 인해 행에 있는 하나의 원소에 액세스하려고 할 때, DRAM 대역폭 활용률이 저하될 수 있습니다. 그 결과, 이미지 라이브러리는 이미지를 메모리에 저장할 때 아래처럼 패딩된 포맷으로 변환합니다.

A padded image format and the concept of pitch.

위 이미지에서 원본 이미지를 3x3이라고 가정하고, 각 DRAM의 버스트가 4픽셀을 포함한다고 가정하겠습니다. 만약 패딩이 없다면 \(M_{1,0}\)은 DRAM 첫 번째 행의 버스트 유닛에 위치하고, \(M_{1,1}\)과 \(M_{1,2}\)는 다음 DRAM 버스트 유닛에 위치하게 됩니다. 따라서 row 1에 액세스하려면 두 개의 DRAM 버스트가 필요하고 이는 메모리 대역폭의 절반이 낭비되는 꼴이 됩니다. 이러한 비효율적인 현상을 해결하기 위해 라이브러리는 각 행 끝에 요소 하나를 패딩합니다. 패딩된 요소를 사용하면 각 행이 이제 전체 DRAM 버스트 유닛을 차지하고, 따라서 1행 또는 2행의 데이터에 액세스할 때 DRAM 버스트 유닛 하나로 전체 행에 액세스할 수 있습니다.

 

패딩을 사용하면 패딩된 요소들에 의해서 이미지 행렬은 확장됩니다. 하지만, 이미지 blur와 같은 연산 중에 패딩된 요소들은 연산에 포함되지 않도록 해야합니다. 따라서 연산을 위해서 이미지의 원래 너비와 높이뿐만 아니라 모든 행의 실제 시작 위치를 제대로 찾을 수 있도록 패딩된 요소에 대한 정보도 제공되어야 합니다. 패딩된 요소의 정보는 위 이미지에서 패딩된 행렬의 pitch로 전달됩니다.

 

아래 이미지는 패딩된 이미지 행렬의 row-major 레이아웃에서 이미지 픽셀 요소들에 액세스하는 방법을 보여줍니다.

Row-major layout of a 2D image matrix with padded elements.

2D 행렬을 1D로 선형화하여 표현하였습니다. \(M_3, M_7, M_11\)은 패딩된 요소로 더미로 취급됩니다. 그리고 선형화된 픽셀 요소의 인덱스를 계산하기 위해서 표현식에 width 대신 pitch가 사용됩니다. 하지만, 원래 요소들만 연산에 사용되도록 루프 바운드는 width를 사용해야 합니다.

 

(pitch를 적용하는 경우에는 DRAM의 버스트 크기를 알아야 하고, 코드도 복잡해져서 pitch는 무시하였습니다.)

이제 Tiled 2D Convolution 커널 함수를 디자인해보겠습니다. 타일의 크기(O_TILE_WIDTH)가 16x16으로 고정하도록 하겠습니다.

#define O_TILE_WIDTH 16
#define MAX_KERNEL_WIDTH 10

__global__
void convolution2D_tiled(float* in, float* out, int width, int height, int channels, int kernel_width)
{
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int tz = threadIdx.z;
    int ch = blockDim.z*blockIdx.z + tz;
    int row_o = blockIdx.y*O_TILE_WIDTH + ty;
    int col_o = blockIdx.x*O_TILE_WIDTH + tx;
    int row_i = row_o - kernel_width/2;
    int col_i = col_o - kernel_width/2;

    __shared__ float in_tile[O_TILE_WIDTH + MAX_KERNEL_WIDTH - 1][O_TILE_WIDTH + MAX_KERNEL_WIDTH - 1];
    if ((row_i >= 0) && (row_i < height) && (col_i >= 0) && (col_i < width)) {
        in_tile[ty][tx] = in[ch*width*height + row_i*width + col_i];
    }
    else {
        in_tile[ty][tx] = 0;
    }

    __syncthreads();

    float val = 0.f;
    if (ty < O_TILE_WIDTH && tx < O_TILE_WIDTH && ch < channels) {
        for (int i = 0; i < kernel_width; i++) {
            for (int j = 0; j < kernel_width; j++) {
                val += in_tile[i+ty][j+tx] * M[i*kernel_width + j];
            }
        }

        if (row_o < height && col_o < width)
            out[ch*width*height + row_o*width + col_o] = val;
    }
}

 

먼저 각 스레드 블록 별로 처리할 입출력 Tile을 설계해야 합니다. (line 7-14)

위 이미지와 같이 Input 타일은 Halo Cell을 포함해야하며, 각 방향의 Halo Cell의 개수만큼 Output 타일에서 확장되어야 합니다. line 7-14의 코드를 살펴보면,

    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int tz = threadIdx.z;
    int ch = blockDim.z*blockIdx.z + tz;
    int row_o = blockIdx.y*O_TILE_WIDTH + ty;
    int col_o = blockIdx.x*O_TILE_WIDTH + tx;
    int row_i = row_o - kernel_width/2;
    int col_i = col_o - kernel_width/2;

먼저 output의 요소의 인덱스를 계산하는데, 위 코드에서 col_o와 row_o에 해당합니다. 그리고 공유 메모리에 적재할 input 요소의 인덱스는 col_i와 row_i에 해당합니다. 

 

인덱스 계산이 완료되면, 이제 공유 메모리에 input 타일을 적재해야합니다. in_tile 배열은 디바이스의 공유 메모리에 위치하며, halo cell을 포함할 수 있을만큼 input 타일의 크기보다 더 큰 사이즈로 할당합니다.

모든 스레드가 공유 메모리 적재에 참여하지만, input 타일 요소의 y및 x 인덱스가 유효한 범위 내에 있는지 체크해주어야 합니다. 유효한 범위 내에 존재하지 않는 인덱스의 경우에는 해당 요소는 0으로 저장해야하고, 이 요소들은 ghost cell이 됩니다. 

    __shared__ float in_tile[O_TILE_WIDTH + MAX_KERNEL_WIDTH - 1][O_TILE_WIDTH + MAX_KERNEL_WIDTH - 1];
    if ((row_i >= 0) && (row_i < height) && (col_i >= 0) && (col_i < width)) {
        in_tile[ty][tx] = in[ch*width*height + row_i*width + col_i];
    }
    else {
        in_tile[ty][tx] = 0;
    }

    __syncthreads();

위 코드는 사실 채널 수가 한 개인 경우에만 동작하는데, 실제 코드에서는 채널 수만큼 그리드의 z차원을 추가할 것이여서 큰 문제가 되지 않습니다. 대신 각 채널에 해당되는 요소에 접근하기 위해 line 3처럼 인덱싱해야 합니다.

 

마지막으로 다음 코드는 공유 메모리의 input 원소들을 사용하여 output을 계산합니다.

    float val = 0.f;
    if (ty < O_TILE_WIDTH && tx < O_TILE_WIDTH && ch < channels) {
        for (int i = 0; i < kernel_width; i++) {
            for (int j = 0; j < kernel_width; j++) {
                val += in_tile[i+ty][j+tx] * M[i*kernel_width + j];
            }
        }

        if (row_o < height && col_o < width)
            out[ch*width*height + row_o*width + col_o] = val;
    }

위에서 input 타일 요소들을 채울 때 input 타일 요소들은 output 요소보다 많다는 것을 볼 수 있습니다. 즉, output 타일의 픽셀 수보다 스레드 블록의 스레드가 더 많습니다(input 타일 요소를 채울 수 있을 만큼의 스레드가 필요하기 때문). 따라서, 위 코드 line 2의 if문은 O_TILE_WIDTH보다 작은 스레드만 output 요소 계산에 참여할 수 있도록 해줍니다. 바로 다음에 이어지는 2중 for문은 컨볼루션 커널과 input 타일의 픽셀 값에 대해 곱셈 및 누적 연산을 수행합니다.

그리고 마지막으로 output 요소가 유효한 범위에 있는 스레드만 결과값을 저장하도록 합니다.

 

전체 코드는 아래 링크에서 참조 가능합니다.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/convolution2D/conv2D.cu

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

 

1D 컨볼루션처럼 간단하게 성능 분석을 진행해보면.. 타일을 사용하지 않은 기본 커널에서 스레드 블록의 모든 스레드는 입력에 대해 (kernel_width)\(^{2}\)의 액세스를 수행합니다. 따라서 각 스레드 블록(O_TILE_WIDTH x O_TILE_WIDTH)은 입력에 대해 총 (kernel_width)\(^{2}\) * (O_TILE_WIDTH)\(^{2}\)의 액세스를 수행합니다.

 

타일을 적용하면 스레드 블록의 모든 스레드는 협력하여 입력을 하나의 타일로 로드합니다. 따라서 스레드 블록에서 입력에 액세스하는 총 횟수는 (O_TILE_WIDTH + kernel_width - 1)\(^{2}\) 입니다.

 

따라서, 기본 2D 컨볼루션과 Tiled 2D 컨볼루션 간의 액세스 비율은 다음과 같습니다.

\[\text{(kernel_width)}^2 * \text{(O_TILE_WIDTH)}^2 / \text{(O_TILE_WIDTH + kernel_width - 1)}^2\]

 

아래 표는 output 타일 크기인 O_TILE_WIDTH에 따른 액세스 비율을 보여주고 있습니다.

O_TILE_WIDTH의 크기가 커질수록 액세스가 크게 감소하는 것을 볼 수 있습니다. 다만, O_TILE_WIDTH가 커질수록 input 타일을 저장하는데 필요한 공유 메모리도 커진다는 것을 감안해야 합니다.

 

실제로 기본 2D 컨볼루션과 Tiled 2D 컨볼루션을 실행하여 성능을 비교해봤습니다. 이전 포스팅의 1D 컨볼루션과 마찬가지로 두 경우에서 큰 성능의 차이는 발견하지 못했습니다.... ㅠ (로직 상 문제가 없어보이는데...)

 

4096x2048x1 입력에 대해

아래의 결과는 기본 2D 컨볼루션의 결과입니다.

그리고 Tiled 2D 컨볼루션의 결과는 다음과 같습니다.