[OpenCL 입문 시리즈 3편] 커널 작성 & 빌드: 기본 문법 따라하기

안녕하세요! 지난 글에서 OpenCL 플랫폼과 디바이스 개념을 이해하고, 원하는 디바이스를 선택하는 방법까지 살펴봤습니다. 이제 본격적으로 커널(Kernel) 작성과 빌드 과정을 알아볼 차례예요. OpenCL 커널은 실제로 GPU나 CPU 등 디바이스에서 병렬 실행되는 코드로, 우리가 가속하고 싶은 연산의 핵심 부분이라고 할 수 있죠.

이번 글에서는 다음 내용을 다룹니다.

• OpenCL 커널 기본 문법 소개
• 호스트 코드에서 커널 빌드 방법 복습 및 확장
• 인자 전달 및 메모리 관리 구조 이해
• CUDA 커널과 비교해 보는 코드 작성 차이점
• 참고할만한 유튜브 자료 링크

앞선 글에서 kernel.cl 파일 안에 간단한 add_one 커널을 작성해봤는데, 이번 글에서는 이보다 조금 더 다양한 형태의 커널 코드를 다뤄보면서 기본 문법을 자연스럽게 익혀보겠습니다.

1. 커널 코드의 기본 구조

OpenCL 커널 코드는 C 언어와 유사한 문법을 사용하고, .cl 파일 안에 작성됩니다. 커널 함수는 __kernel 키워드를 사용해 정의하고, __global 키워드는 메모리를 가리킬 때 사용됩니다. GPU 메모리에 접근할 때는 이와 같은 전용 키워드를 써야 하며, 이는 CUDA에서 __global__, __device__와 같은 키워드를 사용하는 것과 비슷한 맥락입니다.

__kernel void multiply_by_two(__global float* data) {
    int gid = get_global_id(0); 
    data[gid] *= 2.0f;
}

여기서 get_global_id(0)는 현재 워크아이템(work-item)의 ID를 얻는 함수로, OpenCL 런타임이 정해준 전역 인덱스를 반환해 줍니다. CUDA에서 threadIdx.x, blockIdx.x 등을 통해 스레드 인덱스를 얻었던 것과 유사한 역할이에요.

2. 커널 빌드 과정 복습 및 확장

호스트 코드에서 커널 코드를 빌드하고 실행하기 위한 과정은 지난 1편 “Hello OpenCL!” 예제에서 한 번 다뤘었어요. 간단히 복습하자면 다음 단계로 진행됩니다.

1. clCreateProgramWithSource로 .cl 코드를 프로그램 오브젝트로 만듭니다.
2. clBuildProgram으로 해당 프로그램을 디바이스에 맞게 빌드합니다.
3. clCreateKernel으로 빌드된 프로그램에서 특정 커널 함수를 가져옵니다.

이 과정은 CUDA에서 커널 함수를 C++ 코드 안에서 직접 컴파일하고, <<< >>> 연산자로 호출하는 방식과 비교했을 때 다소 번거롭습니다. 하지만 OpenCL은 런타임 빌드 방식을 통해 다양한 디바이스별 최적화를 런타임에 적용할 수 있고, 때에 따라서는 JIT(Just-In-Time) 컴파일 개념을 활용할 수도 있습니다.

3. 더 복잡한 커널 예제: 벡터 덧셈

이제 벡터 두 개를 받아서 결과 벡터를 만들어내는 간단한 연산을 예로 들어보겠습니다. CUDA로 작성한 벡터 덧셈 예제를 본 적 있다면, OpenCL에서도 비슷한 구조라는 걸 알 수 있을 거예요.

kernel.cl

__kernel void vector_add(__global const float* A,
                         __global const float* B,
                         __global float* C,
                         const int n) {
    int gid = get_global_id(0);
    if (gid < n) {
        C[gid] = A[gid] + B[gid];
    }
}

이 커널은 길이 n의 벡터 A와 B를 더해서 C에 저장하는 단순한 연산을 수행합니다. __global const float* A와 같이 const를 붙여서 읽기만 하는 메모리인지 명확히 해둘 수도 있어요.

main.cpp (호스트 코드 스니펫)

이번에는 이 커널을 사용하는 호스트 코드를 간략히 보여드릴게요. 지난 글에서 확인한 단계(프로그램 생성, 빌드, 커널 생성)는 동일하고, 차이점은 커널 인자 설정과 워크 사이즈 지정입니다.

#include <CL/cl.h>
#include <iostream>
#include <vector>
#include <stdexcept>

int main() {
    // (플랫폼, 디바이스 선택, 컨텍스트, 큐 생성 로직은 이전 글 참고)
    // 여기서는 이미 context, queue, device가 있다고 가정
    
    const char* kernelSource = R"CLC(
        __kernel void vector_add(__global const float* A,
                                 __global const float* B,
                                 __global float* C,
                                 const int n) {
            int gid = get_global_id(0);
            if (gid < n) {
                C[gid] = A[gid] + B[gid];
            }
        }
    )CLC";

    cl_int err;
    cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, nullptr, &err);
    if (err != CL_SUCCESS) throw std::runtime_error("프로그램 생성 실패");
    
    err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
    if (err != CL_SUCCESS) {
        // 빌드 에러 로그 출력 (이전 글 참고)
    }

    cl_kernel kernel = clCreateKernel(program, "vector_add", &err);
    if (err != CL_SUCCESS) throw std::runtime_error("커널 생성 실패");

    // 벡터 준비
    const int n = 1000;
    std::vector<float> A(n, 1.0f), B(n, 2.0f), C(n, 0.0f);

    // 버퍼 생성
    cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                 sizeof(float)*n, A.data(), &err);
    cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                 sizeof(float)*n, B.data(), &err);
    cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                                 sizeof(float)*n, nullptr, &err);

    // 커널 인자 설정
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC);
    clSetKernelArg(kernel, 3, sizeof(int), &n);

    // 커널 실행: 글로벌 워크사이즈 = n
    size_t global_work_size = n;
    err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_work_size, nullptr, 0, nullptr, nullptr);
    if (err != CL_SUCCESS) throw std::runtime_error("커널 실행 실패");

    // 결과 읽어오기
    clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(float)*n, C.data(), 0, nullptr, nullptr);

    std::cout << "결과 확인 (일부): " << C[0] << ", " << C[1] << ", ...\n";

    // 자원 해제
    clReleaseMemObject(bufA);
    clReleaseMemObject(bufB);
    clReleaseMemObject(bufC);
    clReleaseKernel(kernel);
    clReleaseProgram(program);

    return 0;
}

위 코드를 실행하면, C 벡터의 원소가 모두 3.0f(1.0 + 2.0)가 되어 있어야 합니다.

4. CUDA 커널과 비교하기

CUDA 커널 코드와 비교하면 OpenCL 커널은 다음과 같은 특징이 있어요.

• 런타임 빌드: CUDA는 NVCC를 통해 컴파일된 PTX/바이너리를 사용하지만, OpenCL은 기본적으로 소스 코드를 런타임에 빌드합니다.
• 인자 설정 방식: CUDA는 <<< >>>로 런치하면서 런치 파라미터를 지정하고, 커널 내에 직접 함수를 호출하는 느낌이라면, OpenCL은 clSetKernelArg로 각 인자를 하나씩 설정하는 절차가 필요해요.
• 메모리 관리: CUDA에서는 cudaMalloc, cudaMemcpy를 사용하고, OpenCL에서는 clCreateBuffer, clEnqueueWriteBuffer를 사용합니다. 함수 이름만 다를 뿐 개념적으로 유사한 부분이 많습니다.

둘 다 병렬 실행 모델을 다루지만, CUDA가 NVIDIA만을 위한 깔끔한 통합 환경이라면, OpenCL은 범용 디바이스 지원을 위해 좀 더 많은 수동 설정을 요구하는 구조입니다.

5. 마무리

이번 글에서는 커널 작성 기본 문법과 빌드 과정을 좀 더 자세히 알아봤습니다. 이제 OpenCL 커널 코드를 볼 때 __kernel, __global, get_global_id 같은 키워드가 조금 더 익숙하게 느껴지리라 생각해요. 다음 글에서는 메모리 모델과 버퍼, 이미지, 파라미터를 좀 더 체계적으로 다루며, OpenCL의 메모리 관리 전략을 파헤쳐보겠습니다.

유용한 링크 & 리소스

• OpenCL 공식 홈페이지 : 커널 프로그래밍 레퍼런스
• Khronos Group GitHub : 예제 및 샘플 코드
• C++ Reference : 최신 C++ 문법 참고