Array of Structures 와 Structure of Arrays

References

• Professional CUDA C Programming

Contents

• Array of Structures vs. Structure of Arrays

---

여기 두 가지 데이터를 구성하는 방법이 있습니다.

• Array of Structures (AoS)
• Structure of Arrays (SoA)

이 두 가지 방법은 각각 두 가지의 데이터 타입(구조체와 배열)을 활용할 수 있는 다른 방법을 보여주기 때문에 항상 흥미로운 주제입니다.

 

두 개의 float 데이터를 저장하는 예제를 살펴보겠습니다.

먼저, AoS 방법을 사용하면, innerStruct라는 배열을 정의할 수 있습니다.

struct innerStruct {
	float x;
    float y;
};

그리고 이 구조체의 배열을 정의합니다.

struct innerStruct myAoS[N];

이것이 AoS로 데이터를 구성하는 방법입니다. 관련 데이터를 공간적으로 가깝게 저장하므로, CPU에서 캐시 지역성이 우수합니다. 

 

다음은 SoA 방법입니다.

struct innerArray {
	float x[N];
    float y[N];
};

이는 원래 구조의 각 필드에 대한 모든 값이 그 필드의 배열로 구분됩니다. 이렇게 하면 인접한 포인트의 데이터가 함께 저장되지만, 하나의 포인트에 대한 필드 값들이 여러 배열들에 나누어져서 저장됩니다. 위 구조체는 다음과 같이 변수를 정의할 수 있습니다.

struct innerArray mySoA;

 

아래 그림은 AoS와 SoA 방법에서의 메모리 레이아웃을 보여줍니다.

GPU에서 AoS 포맷의 데이터를 저장하고, x 필드에만 필요한 작업을 수행하면 y의 값도 암시적으로 32바이트의 segment 또는 128바이트의 캐시라인에 로드되므로 50%의 bandwidth 손실이 발생합니다. 또한 AoS 포맷은 불필요한 y의 값 때문에 L2 캐시 공간을 낭비합니다.

 

SoA 방식으로 데이터를 저장하면 GPU 메모리 bandwidth를 최대한으로 활용할 수 있습니다. 동일한 필드의 요소들이 연속적으로 위치하기 때문에 GPU에서 SoA 레이아웃은 coalesced memory access를 제공하며 더 효율적으로 global memory 활용할 수 있도록 합니다.

 

AOS versus SOA
많은 병렬 프로그래밍 패러다임에서, 특히 SIMD 스타일에서 SOA를 더 선호합니다. CUDA C 프로그래밍에서도 SoA가 선호되는데, 이는 데이터 요소들이 효율적으로 global memory에 coalesced access할 수 있도록 정렬되기 때문입니다.

 

각 데이터 타입에 따라서 성능에 어떠한 영향을 미치는지 직접 확인해보도록 하겠습니다.

 

---

Example for AoS Data layout

다음은 AoS 레이아웃을 사용하여 구현된 커널입니다. 

__global__
void testInnerStruct(innerStruct* data, innerStruct* result, const int N)
{
    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx < N) {
        innerStruct tmp = data[idx];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[idx] = tmp;
    }
}

간단하게, x 필드의 값에 10.0을 더하고 y 필드의 값에는 20.0을 더해서 결과에 저장합니다.

 

전체 코드는 아래 링크를 참조해주세요.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/AOSandSOA/AoS.cu

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

 

입력의 크기는 \(2^20\)으로 고정시키고, block 사이즈를 (128, 1)로 설정하고 컴파일 후 실행해보겠습니다.

다음의 커맨드로 컴파일하고,

nvcc -O3 -o AoS AoS.cu -I..

실행하면,

위와 같은 출력 결과를 확인할 수 있습니다.

 

Nsight System으로 global load 효율성과 global store 효율성을 프로파일링해보겠습니다.

다음의 커맨드로 측정하고 싶은 것만 설정하여서 프로파일링할 수 있습니다.,

ncu.bat --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct,smsp__sass_average_data_bytes_per_sector_mem_global_op_st.pct ./AoS.exe

다음의 결과를 얻을 수 있는데, 첫 번째 항목이 global load 효율성이고 두 번째 항목이 global store 효율성입니다.

load와 store 모두 50%의 효율성을 보여주고 있습니다. 이는 x와 y의 값이 메모리에 인접하기 때문에 memory transaction이 특정 필드의 값을 로드하기 위해 수행될 때마다 로드된 바이트의 절반은 다른 필드에 속하기 때문입니다. 따라서 대역폭의 50%가 사용되지 않습니다.

 

Example of SoA Data layout

이번에는 SoA 데이터를 가지고 동일한 커널을 구현하여 테스트해보겠습니다.

사용할 커널 함수는 다음과 같습니다.

__global__
void testInnerArray(innerArray* data, innerArray* result, const int N)
{
    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx < N) {
        float tmpX = data->x[idx];
        float tmpY = data->y[idx];

        tmpX += 10.f;
        tmpY += 20.f;
        result->x[idx] = tmpX;
        result->y[idx] = tmpY;
    }
}

 

마찬가지로 동일한 입력 크기와 동일한 스레드 개수로 프로파일링 하도록 하겠습니다.

전체 코드는 아래 링크를 참조해주세요 !

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/AOSandSOA/SoA.cu

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

 

다음의 커맨드로 컴파일하고,

nvcc -O3 -o SoA SoA.cu -I..

실행하면,

위와 같은 결과를 볼 수 있습니다.

 

아까 전과 동일하게 global load/store 효율성을 측정해보도록 하겠습니다.

다음의 커맨드로 실행합니다.

ncu.bat --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct,smsp__sass_average_data_bytes_per_sector_mem_global_op_st.pct ./SoA.exe

global load/store 효율성이 모두 100%로 측정됩니다. 하나의 memory transaction에 의해서 각 액세스가 처리되고 있다는 것을 알 수 있습니다.

 

 

이처럼 병렬 프로그래밍에서는 SoA가 AoS보다 실행/효율성 측면에서 더 유리하다는 것을 알 수 있습니다.. !