[SYCL 입문 #5] 매트릭스 곱 실용 예제와 성능 개선 아이디어

SYCL 입문 시리즈의 다섯 번째 글입니다. 이전 글(#4)에서 ND-Range 활용, USM(United Shared Memory) 소개, 로컬 액세서 등 성능 최적화의 기초 개념을 다뤄보았습니다. 이제는 지금까지 습득한 개념을 종합해 조금 더 실용적인 예제를 만들어보겠습니다. 이번 글에서는 매트릭스 곱(Matrix Multiplication) 연산을 SYCL로 구현하고, ND-Range 설정, 로컬 메모리(로컬 액세서), USM 혹은 버퍼 기반 접근을 통해 성능을 개선하는 아이디어를 제시합니다.

목표

• 매트릭스 곱(행렬 A( MxK ) × B( KxN ) = C( MxN ))를 SYCL 커널로 구현
• ND-Range 설정을 통해 워크그룹 크기, 워크아이템 분배 조정
• 로컬 액세서(local_accessor) 활용해 매트릭스 타일링(tiling)으로 메모리 접근 최적화 개념 맛보기
• USM 또는 버퍼 기반 접근 중 하나 선택해 구현 (초심자는 버퍼 기반부터 시작)
• CMake 빌드 시 백엔드 변경(NVIDIA, Qualcomm GPU) 응용 가능

매트릭스 곱 개념 복습

매트릭스 곱 C = A×B에서 C[i,j] = Σ_k A[i,k]*B[k,j]. 이를 병렬로 처리하려면 각 C[i,j]를 독립적으로 계산할 수 있습니다. SYCL 커널에서 (i,j)에 해당하는 워크아이템이 A의 i행과 B의 j열을 순회하면서 곱하고 더하는 과정을 거칩니다.

기본 구현(로컬 메모리 없이) 예:

cgh.parallel_for<class matmul_basic>(
    sycl::nd_range<2>(sycl::range<2>(M,N), sycl::range<2>(16,16)),
    [=](sycl::nd_item<2> item) {
        size_t i = item.get_global_id(0);
        size_t j = item.get_global_id(1);
        float sum = 0.0f;
        for (size_t k=0; k<K; k++) {
            sum += A[i*K + k] * B[k*N + j];
        }
        C[i*N + j] = sum;
    }
);

위 코드는 단순하고 이해하기 쉽지만, 전역 메모리에 A, B에서 매번 k 루프마다 접근하므로 메모리 대역폭 소모가 큽니다.

로컬 액세서를 통한 타일링 기법 도입(기본 아이디어)

고급 최적화: 로컬 메모리(로컬 액세서)를 사용해 A와 B의 일부를 타일(tile) 단위로 워크그룹에 로드한 뒤, 워크그룹 내에서 공유하면서 k 루프를 돌면 전역 메모리 접근 횟수를 줄일 수 있습니다. 이는 CUDA나 OpenCL 최적화와 유사한 패턴입니다.

간단한 예(개념):

1. 워크그룹당 (16x16) 타일을 A와 B에서 로컬 메모리에 로드
2. k 루프마다 작은 타일 크기 단위로 곱하고 더하기
3. k 루프 모두 완료 후 C[i,j]에 결과 저장

이 패턴은 이 글에서는 개념만 설명하고, 실제 코드는 간단 버전(로컬 메모리 없이)과 로컬 메모리 버전 두 가지로 비교해 볼 수 있습니다.

예제 코드 (단순 버전)

먼저 로컬 메모리 없이 버퍼 기반으로 구현한 매트릭스 곱 예제를 작성하겠습니다. ND-Range로 (16,16) 워크그룹을 사용해보겠습니다.

디렉토리 구조

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

main.cpp (단순 버전, 로컬 메모리 미사용)

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

int main() {
    try {
        sycl::queue q;
        std::cout << "Running on: " << q.get_device().get_info<sycl::info::device::name>() << "\n";

        size_t M=256, K=256, N=256;
        std::vector<float> A(M*K, 1.0f), B(K*N, 1.0f), C(M*N, 0.0f);

        {
            sycl::buffer<float,1> bufA(A.data(), sycl::range<1>(M*K));
            sycl::buffer<float,1> bufB(B.data(), sycl::range<1>(K*N));
            sycl::buffer<float,1> bufC(C.data(), sycl::range<1>(M*N));

            auto start = std::chrono::high_resolution_clock::now();
            q.submit([&](sycl::handler& cgh) {
                auto aacc = bufA.get_access<sycl::access::mode::read>(cgh);
                auto bacc = bufB.get_access<sycl::access::mode::read>(cgh);
                auto cacc = bufC.get_access<sycl::access::mode::write>(cgh);

                cgh.parallel_for<class matmul_basic>(
                    sycl::nd_range<2>(sycl::range<2>(M,N), sycl::range<2>(16,16)),
                    [=](sycl::nd_item<2> item) {
                        size_t i = item.get_global_id(0);
                        size_t j = item.get_global_id(1);
                        float sum = 0.0f;
                        for (size_t k=0; k<K; k++) {
                            sum += aacc[i*K + k] * bacc[k*N + j];
                        }
                        cacc[i*N + j] = sum;
                    }
                );
            });
            q.wait();
            auto end = std::chrono::high_resolution_clock::now();
            double elapsed = std::chrono::duration<double, std::milli>(end-start).count();
            std::cout << "Matmul done in " << elapsed << " ms\n";
        } // 스코프 끝나면 C에 결과 반영

        std::cout << "C[0,0] = " << C[0] << "\n"; // 모두 256.0f (1*1*256)
    } 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(matmul_example CXX)

set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)

if(NOT DEFINED SYCL_COMPILER)
    set(SYCL_COMPILER "dpcpp")
endif()

set(CMAKE_CXX_COMPILER ${SYCL_COMPILER})

add_executable(matmul_example src/main.cpp)

빌드 & 실행:

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

출력 예시:

Running on: NVIDIA GPU ...
Matmul done in xxx ms
C[0,0] = 256

NVIDIA나 Qualcomm GPU에서도 이전 글에서 설명한 백엔드 선택 및 device_selector를 응용하면 동일 코드가 실행될 수 있습니다.

로컬 메모리 사용(개념 소개)

로컬 메모리를 사용하면 다음과 같은 식으로 바꿀 수 있습니다. (코드 전체는 여기서 다 제시하지 않고 개념만 설명)

1. ND-Range를 (M,N) 범위로 설정하되, 워크그룹 크기를 (16,16)으로 설정
2. sycl::local_accessor<float,2>를 이용해 A와 B의 부분 타일을 로컬 메모리에 로드
3. 각 워크그룹 내에서 k 루프를 돌 때마다 타일 단위로 로컬 메모리에 적재한 데이터를 활용, 전역 메모리 접근 최소화

이런 최적화는 성능을 크게 향상시킬 수 있습니다. 다만, 입문 단계에서는 개념만 알고 있다가 향후 성능 개선 필요 시 적용해보면 됩니다.

USM으로 구현 시 장단점

USM을 사용하면 malloc_shared나 malloc_device로 메모리를 할당하고, 포인터 기반으로 접근할 수 있습니다. 이 경우 버퍼/액세서 대신 단순 포인터 활용:

float* A_data = sycl::malloc_shared<float>(M*K, q);
float* B_data = sycl::malloc_shared<float>(K*N, q);
float* C_data = sycl::malloc_shared<float>(M*N, q);

// 초기화 후 커널에서 data[idx] 형태로 바로 접근 가능

USM은 단순하지만, 메모리 전송을 런타임에 맡기지 않고 프로그래머가 더 관리에 신경 써야 할 수 있습니다. 입문자는 버퍼 기반 모델로 시작한 뒤, 필요 시 USM으로 전환하는 전략을 권장합니다.

정리 및 다음 글 예고

이번 글에서 ND-Range를 통한 워크그룹 크기 조정, 로컬 액세서 기반 타일링 개념 소개, USM을 통한 포인터 기반 메모리 관리 등 성능 및 최적화 방향성을 제시했습니다. 입문 단계에서는 모든 최적화를 당장 적용할 필요는 없지만, SYCL이 이런 가능성을 제공한다는 점을 인지하고 있으면, 향후 더 복잡한 프로젝트에서 유용하게 활용할 수 있습니다.

다음 글(#6)에서는 지금까지 배운 개념을 바탕으로 더 큰 규모의 예제나, 간단한 벤치마크를 통해 성능 차이를 관찰하는 방법을 다루며, 디버깅/프로파일링 툴과 결합하는 기초적인 성능 분석 과정을 살펴보겠습니다.