[SYCL 입문 #4] ND-Range 활용과 메모리 최적화 기법 소개

SYCL 입문 시리즈의 네 번째 글입니다. 지난 글(#3)에서 SYCL의 메모리 모델과 커널 작성 패턴, 워크아이템/워크그룹 개념을 다뤄봤습니다. 이제 한 걸음 더 나아가, SYCL에서 ND-Range(nd_range) 구성을 활용하고, 더 복잡한 메모리 관리 기법(예: 다양한 메모리 속성, USM(United Shared Memory) 소개) 등 성능 최적화 전략의 기초를 살펴보겠습니다. 이 글은 여전히 입문자를 대상으로 하며, 모든 단계를 가능한 한 자세히 설명하므로, 독자가 직접 따라 해볼 수 있습니다.

다만, USM 등 일부 개념은 SYCL 1.2.1 표준 이후 확장된 기능으로, 구현체(예: oneAPI DPC++, hipSYCL)마다 지원 상황이 다를 수 있으므로, 실습 전에 구현체 문서를 참고하시기 바랍니다.

목표

• ND-Range(nd_range) 사용법: 워크그룹 크기 지정과 커널 런치 최적화
• 메모리 접근 최적화 개념: 호스트 메모리 접근 최소화, 디바이스 로컬 메모리(로컬 액세서) 사용
• USM(United Shared Memory) 개념 소개: 버퍼 대신 포인터 기반 메모리 관리로 더 직접적인 메모리 제어
• CMake 빌드 시 특정 옵션을 통해 USM 지원 여부 확인

ND-Range(nd_range)란?

이전 글에서는 parallel_for(range<1>) 또는 parallel_for(range<2>) 형태로 워크아이템 범위를 지정했습니다. 여기서 워크그룹 크기를 세밀하게 조정하려면 nd_range를 사용합니다.

예시:

cgh.parallel_for<class kernel_ndrange>(
    sycl::nd_range<2>(sycl::range<2>(1024,1024), sycl::range<2>(16,16)),
    [=](sycl::nd_item<2> item) {
        auto i = item.get_global_id(0);
        auto j = item.get_global_id(1);
        // (i,j)에 대한 연산
    }
);

여기서 (1024,1024)는 전체 워크아이템 수, (16,16)는 워크그룹 크기입니다. 이를 통해 GPU의 워프/웨이브프런트 크기와 맞춰 최적화를 시도할 수 있습니다. NVIDIA GPU에서 32스레드 단위로 묶을 때나 특정 워크그룹 크기를 고정해 메모리 접근 패턴을 최적화하는데 유용합니다.

로컬 액세서(로컬 메모리) 활용

SYCL에서는 로컬 메모리를 액세서를 통해 사용할 수 있습니다. 로컬 메모리는 워크그룹 단위로 공유되는 메모리로, CUDA의 shared memory와 유사한 개념입니다.

예:

cgh.parallel_for<class kernel_local>(
    sycl::nd_range<2>(sycl::range<2>(1024,1024), sycl::range<2>(16,16)),
    [=](sycl::nd_item<2> item) {
        // 로컬 액세서 선언
        // 예: float 타입의 16x16 로컬 메모리
        // 이 로컬 액세서는 command group 밖에서 선언해 cgh에 전달해야 함
    }
);

로컬 액세서 사용 시 다음과 같은 단계 필요:

1. sycl::local_accessor를 핸들러(cgh) 안에서 생성
2. 커널 내부에서 nd_item을 통해 로컬 메모리에 접근, 워크그룹 내 연산 최적화
3. 이로써 전역 메모리에 대한 접근 횟수를 줄여 성능 향상 가능

로컬 액세서는 초기 입문 단계에서는 Optional이지만, 큰 데이터 처리 시 성능 최적화에 도움을 줍니다.

USM(United Shared Memory) 소개

기존 SYCL 버퍼/액세서 모델에서는 메모리 관리를 SYCL 런타임에 위임했는데, USM은 포인터 기반 메모리 관리로 개발자가 더 직접적으로 메모리를 컨트롤할 수 있습니다. CUDA malloc, memcpy에 더 가까운 모델입니다.

• USM 할당: sycl::malloc_shared, malloc_host, malloc_device 함수를 통해 포인터를 할당
• 직접 memcpy: sycl::memcpy 호출 가능
• 호스트/디바이스 간 메모리 전송을 개발자가 직접 제어
• 초기 입문자에게는 조금 복잡할 수 있지만, 향후 성능 최적화나 특정 메모리 패턴 구현 시 유용

예제(개념 예시):

float* data = sycl::malloc_shared<float>(N, q);
for (int i=0; i<N; i++) data[i]=1.0f;
q.parallel_for(sycl::range<1>(N), [=](sycl::id<1> i) {
    data[i]*=2.0f;
});
q.wait();
std::cout << "data[0]=" << data[0] << "\n";
sycl::free(data, q);

USM은 구현체마다 지원 수준이 다를 수 있으므로, oneAPI DPC++는 USM 지원이 우수합니다.

빌드 시 옵션 확인

CMake 빌드 시 -D BACKEND=cuda나 -D BACKEND=opencl 등을 통해 백엔드를 결정했던 것처럼, USM 사용 시에도 특정 컴파일러 옵션이 필요할 수 있습니다. 예를 들어 oneAPI DPC++에서는 USM 기본 지원이 있고, hipSYCL에서도 특정 플래그로 USM 지원을 활성화할 수 있습니다. 구현체 문서 확인 후 CMake에 -DUSE_USM=ON 등 옵션을 추가해 조건부로 USM 코드를 컴파일하도록 할 수 있습니다.

예제 코드 (ND-Range & USM 맛보기)

아래는 간단한 예제입니다. N=1024 길이의 배열을 USM으로 할당하고, nd_range로 디스패치해 각 원소를 5배로 만드는 예제입니다. (NVIDIA GPU나 CPU, OpenCL 백엔드 등 상황에 맞게 백엔드를 바꿔 시도할 수 있습니다.)

#include <CL/sycl.hpp>
#include <iostream>

int main() {
    sycl::queue q; // 기본 디바이스 선택
    std::cout << "Running on: " << q.get_device().get_info<sycl::info::device::name>() << "\n";

    size_t N=1024;
    // USM 할당 (shared)
    float* data = sycl::malloc_shared<float>(N, q);
    for (size_t i=0; i<N; i++) data[i]=1.0f;

    q.parallel_for(
        sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(64)), 
        [=](sycl::nd_item<1> item) {
            size_t idx = item.get_global_id(0);
            data[idx] *= 5.0f;
        }
    );
    q.wait();

    std::cout << "data[0] = " << data[0] << ", data[100] = " << data[100] << "\n";
    std::cout << "All elements should now be 5.0\n";

    sycl::free(data, q); // USM free
    return 0;
}

빌드 후 실행 시:

Running on: NVIDIA GPU ... (또는 Intel GPU, CPU 등)
data[0] = 5, data[100] = 5
All elements should now be 5.0

ND-Range로 64개 워크아이템씩 그룹핑했고, USM으로 메모리를 직접 할당/해제했습니다. 이는 입문자가 바로 이 방식을 쓸 필요는 없지만, 향후 최적화나 벤더별 메모리 특성 활용 시 유용한 기법입니다.

정리 및 다음 글 예고

이번 글에서 ND-Range를 활용해 워크그룹 크기를 제어하고, 로컬 액세서나 USM 등 고급 메모리 접근 기법을 소개했습니다. 이러한 기능을 익히면 SYCL 코드 성능 최적화와 다양한 하드웨어 특성 활용이 용이해집니다.

다음 글(#5)에서는 SYCL을 활용한 조금 더 실용적인 예제, 예를 들어 매트릭스 곱 연산 등을 다루어, 지금까지 배운 메모리 모델, 커널 작성 패턴, ND-Range 구성을 종합적으로 적용해보겠습니다. 또한 간단한 성능 개선 아이디어도 제시해볼 예정입니다.