[SYCL 입문 #3] 메모리 모델 & 커널 작성 패턴 이해하기

SYCL 입문 시리즈의 세 번째 글입니다. 이전 글(#2)에서 NVIDIA GPU나 Qualcomm GPU를 대상으로 SYCL 코드를 실행하기 위한 백엔드 선택, 디바이스 셀렉터 사용 방법을 살펴봤습니다. 이제 한 단계 더 나아가, SYCL의 메모리 모델과 커널 작성 패턴을 자세히 다루겠습니다. 이 글은 여전히 입문자를 대상으로 하며, 가능한 한 모든 단계를 구체적으로 설명하여, 독자가 직접 따라 할 수 있도록 합니다.

SYCL에서 메모리와 커널(디바이스 코드) 작성은 SYCL이 제공하는 추상화 덕분에 C++ 람다 표기법을 활용하고, 버퍼(buffer)와 액세서(accessor)를 통해 호스트-디바이스 데이터를 쉽게 관리할 수 있습니다. 또한 워크아이템(work-item), 워크그룹(work-group), ND-Range(nd_range) 개념을 간단한 예제를 통해 이해해보겠습니다.

목표

• SYCL 메모리 모델: 버퍼(buffer), 액세서(accessor), 호스트-디바이스 메모리 전송 자동 관리 이해
• 람다 함수를 이용한 SYCL 커널 작성: C++ 스타일로 GPU 커널 구현
• 워크아이템(Work-item), 워크그룹(Work-group), ND-Range 개념 소개
• 간단한 2D 데이터 처리 예제를 통해 SYCL 커널 패턴 실습
• CMake 빌드 & 이전 글에서 사용한 디바이스 선택/백엔드 설정 응용

SYCL 메모리 모델 개념

SYCL은 프로그래머가 명시적으로 memcpy 같은 것을 호출하지 않아도, 버퍼(buffer) 객체를 통해 호스트 메모리를 디바이스 메모리로 자동 전송하고, 커널 실행 후 결과를 다시 호스트 메모리에 반영합니다.

핵심 개념:

• 버퍼(buffer): 호스트 메모리에 있는 STL 컨테이너(예: std::vector)를 SYCL 범위(range)와 연계해 디바이스 메모리와 동기화하는 추상화 객체.
• 액세서(accessor): 커널 내부에서 디바이스 메모리에 접근하기 위한 핸들. get_access<mode>(cgh)를 통해 읽기, 쓰기 모드를 지정하면 SYCL 런타임이 메모리 전송을 관리.
• 스코프(scope)를 벗어날 때 버퍼의 데이터가 자동으로 호스트로 반영.

이로써 개발자는 데이터 전송 로직을 대부분 신경 쓰지 않고, 단지 버퍼와 액세서를 선언해주면 됩니다.

커널 작성: C++ 람다 이용

SYCL에서 커널은 C++ 람다 표현식으로 작성할 수 있습니다.

예:

cgh.parallel_for<class simple_kernel>(sycl::range<1>(N), [=](sycl::id<1> i) {
    acc[i] = acc[i] * 2.0f;
});

• parallel_for에 람다를 전달하면 SYCL 컴파일러가 이를 디바이스 코드로 컴파일
• class simple_kernel은 커널 식별자 역할(고유 클래스명 필요)
• range<1>(N)은 1차원 범위, i는 인덱스
• 2D나 3D 데이터처리를 위해 range<2>나 range<3> 사용 가능

워크아이템 & 워크그룹

• 워크아이템(work-item): 병렬로 실행되는 단위(예: i 인덱스)
• 워크그룹(work-group): 워크아이템들을 그룹으로 묶은 단위
• parallel_for는 ND-Range(nd_range) 개념을 통해 워크그룹 크기와 전체 범위를 지정할 수 있음.

예:

cgh.parallel_for<class kernel2d>(
    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);
        // 2D 인덱스 (i,j)에 대한 연산
    }
);

이 예제에서는 1024x1024 워크아이템을 16x16 워크그룹으로 나누어 실행.

예제 코드 (2D 데이터 처리)

아래는 2D 배열(1024x1024)을 SYCL 커널로 처리하는 간단한 예제입니다. 모든 요소를 3배로 만들고, 결과를 호스트에 반영한 뒤 출력합니다.

디렉토리 구조

my_sycl_kernel/
 ├─ CMakeLists.txt
 └─ src/
    └─ main.cpp

main.cpp 코드

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

int main() {
    try {
        // 이전 글에서 설명한 device_selector나 백엔드 설정을 여기서 응용 가능
        // 일단 기본 device_selector 사용
        sycl::queue q;
        std::cout << "Running on: " << q.get_device().get_info<sycl::info::device::name>() << "\n";

        const size_t N = 1024;
        std::vector<float> data(N*N, 1.0f); // 1024x1024 요소, 모두 1.0

        {
            sycl::buffer<float,2> buf(data.data(), sycl::range<2>(N, N));

            q.submit([&](sycl::handler& cgh) {
                auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
                cgh.parallel_for<class multiply_by_three>(
                    sycl::range<2>(N, N),
                    [=](sycl::id<2> idx) {
                        acc[idx] *= 3.0f;
                    }
                );
            });
            // 커맨드 그룹 범위 끝나면 데이터가 호스트로 반영
        }

        // 결과 확인
        std::cout << "data[0,0] = " << data[0] << ", data[100,100] = " << data[100*N+100] << "\n";
        std::cout << "All elements should now be 3.0\n";
    } catch(sycl::exception const &e) {
        std::cerr << "SYCL Exception: " << e.what() << "\n";
        return 1;
    }

    return 0;
}

CMakeLists.txt 예제

cmake_minimum_required(VERSION 3.10)
project(kernel_example CXX)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)

if(NOT DEFINED SYCL_COMPILER)
    set(SYCL_COMPILER "dpcpp") # 예: oneAPI DPC++
endif()

set(CMAKE_CXX_COMPILER ${SYCL_COMPILER})

add_executable(kernel_example src/main.cpp)

빌드 & 실행:

mkdir build
cd build
cmake -D SYCL_COMPILER=dpcpp ..
make
./kernel_example

출력 예시:

Running on: Intel(R) UHD Graphics [0x3ea0]
data[0,0] = 3, data[100,100] = 3
All elements should now be 3.0

NVIDIA나 Qualcomm GPU를 활용하려면 이전 글(#2)에서 설명한 백엔드 옵션, device_selector를 변경한 뒤 빌드하면 동일 코드가 해당 GPU에서 실행될 수 있습니다.

정리 및 다음 글 예고

이번 글에서는 SYCL 메모리 모델(버퍼, 액세서)과 커널 작성 패턴(C++ 람다 이용), 워크아이템/워크그룹 개념 등을 다루어, SYCL 코드가 단순한 1D 연산을 넘어 2D 데이터 처리까지 확장 가능함을 보았습니다. 이제 다양한 형태의 데이터를 SYCL을 통해 간단하게 처리할 수 있음을 이해했으리라 생각합니다.

다음 글(#4)에서는 SYCL에서 더 복잡한 메모리 관리 기법이나 ND-Range 구성을 다루며, 성능 최적화와 백엔드별 최적화 전략을 살짝 엿보는 과정을 진행할 예정입니다. 또한 호스트 코드에서 디바이스 코드로 데이터를 넘길 때 고려해야 할 점들을 좀 더 깊게 파고들 예정입니다.