[OpenCL 입문 시리즈 6편] 간단한 이미지 처리 예제: 그레이스케일 변환하기

안녕하세요! 지난 글에서는 C++20/23 문법과 OpenCL C++ Wrapper를 활용해 OpenCL 코드를 훨씬 깔끔하게 만드는 방법을 살펴봤습니다. 이번 글에서는 실제로 유용할 수 있는 간단한 이미지 처리 예제를 통해 OpenCL의 이미지 기능을 체험해보려 합니다. 기본적인 버퍼 연산만 다뤘던 것에서 한 단계 나아가, GPU 가속을 활용한 이미지 변환(예: 컬러 이미지를 그레이스케일로 변환)을 시도해보죠.

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

• 이미지 객체를 사용한 데이터 처리
• 호스트 측에서 이미지 로딩(예: stb_image 사용)
• OpenCL 커널에서 read_imagef, write_imagef 함수 활용
• 결과 확인을 위한 이미지 저장
• CUDA와 비교: CUDA 텍스처/서피스와 유사한 개념

1. 준비사항

이미지 처리를 위해서는 입력 이미지를 로딩하고, 결과를 저장하는 과정이 필요합니다. 여기서는 외부 라이브러리인 stb_image와 stb_image_write 같은 헤더 온리 라이브러리를 활용할 수 있어요. 이 라이브러리는 추가 빌드가 필요 없고, 간단한 인터페이스를 제공하므로 예제용으로 좋습니다.

• stb_image.h / stb_image_write.h 다운로드 후 프로젝트에 포함
• 컬러 이미지(.png, .jpg 등)를 준비

2. 이미지 객체 생성하기

이전 글에서 이미지 객체를 다루는 법을 간략히 언급했는데, 이미지 객체를 생성할 때는 clCreateImage (C API) 또는 C++ Wrapper를 통해 cl::Image2D를 사용할 수 있습니다. 이미지의 포맷(CL_RGBA, CL_UNORM_INT8 등)과 크기(width, height), 접근 권한을 지정한 뒤 컨텍스트를 기반으로 이미지를 만듭니다.

예: cl::Image2D 생성

cl::ImageFormat format(CL_RGBA, CL_UNORM_INT8);
cl::Image2D inputImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
                       format, width, height, 0, inputPixels.data());

cl::Image2D outputImage(context, CL_MEM_WRITE_ONLY, 
                        format, width, height);

여기서 inputPixels는 stbi_load를 통해 읽어온 픽셀 데이터(예: std::vector<unsigned char> 형태)라고 가정합니다.

3. 커널 작성: 컬러 → 그레이스케일 변환

그레이스케일 변환은 RGB 값을 합쳐 평균을 내거나, 가중합을 사용해서 더 실제적인 밝기 변환을 할 수 있습니다. 간단히 gray = (R + G + B) / 3.0f를 적용해보죠.

kernel.cl

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

__kernel void to_grayscale(read_only image2d_t srcImage,
                           write_only image2d_t dstImage,
                           int width,
                           int height) {
    int x = get_global_id(0);
    int y = get_global_id(1);
    if (x >= width || y >= height) return;

    float4 pixel = read_imagef(srcImage, sampler, (int2)(x,y));
    float gray = (pixel.x + pixel.y + pixel.z) / 3.0f;
    float4 grayPixel = (float4)(gray, gray, gray, pixel.w);

    write_imagef(dstImage, (int2)(x,y), grayPixel);
}

read_only image2d_t는 읽기 전용 이미지, write_only image2d_t는 쓰기 전용 이미지를 의미합니다. read_imagef와 write_imagef 함수로 픽셀 접근이 가능합니다.

 

CUDA 텍스처 접근과 유사하지만, OpenCL에서는 이미지를 이렇게 전용 함수로 다룬다는 점이 특징입니다.

4. 호스트 코드: 커널 실행

호스트 코드에서는 입력 이미지를 읽고, cl::Image2D로 만들고, 커널을 실행한 뒤 결과 이미지를 cl::CommandQueue::enqueueReadImage로 받아와서 파일로 저장하면 됩니다.

#include <CL/cl2.hpp>
#include "stb_image.h"
#include "stb_image_write.h"
#include <iostream>
#include <vector>

int main() {
    try {
        // (플랫폼, 디바이스, 컨텍스트, 큐 생성은 이전 글 참고)
        // 여기서는 이미 context, queue, device 있다고 가정

        // 이미지 로딩
        int width, height, channels;
        unsigned char* data = stbi_load("input.png", &width, &height, &channels, 4);
        if (!data) {
            std::cerr << "이미지 로딩 실패\n";
            return 1;
        }

        std::vector<unsigned char> inputPixels(data, data + width*height*4);
        stbi_image_free(data);

        cl::ImageFormat format(CL_RGBA, CL_UNORM_INT8);
        cl::Image2D inputImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
                               format, width, height, 0, inputPixels.data());

        cl::Image2D outputImage(context, CL_MEM_WRITE_ONLY, format, width, height);

        // 커널 로드 & 빌드
        const char* kernelSrc = R"CLC(
            __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
            __kernel void to_grayscale(read_only image2d_t srcImage,
                                       write_only image2d_t dstImage,
                                       int width,
                                       int height) {
                int x = get_global_id(0);
                int y = get_global_id(1);
                if (x >= width || y >= height) return;

                float4 pixel = read_imagef(srcImage, sampler, (int2)(x,y));
                float gray = (pixel.x + pixel.y + pixel.z) / 3.0f;
                float4 grayPixel = (float4)(gray, gray, gray, pixel.w);
                write_imagef(dstImage, (int2)(x,y), grayPixel);
            }
        )CLC";

        cl::Program program(context, kernelSrc);
        program.build({device});
        cl::Kernel kernel(program, "to_grayscale");
        kernel.setArg(0, inputImage);
        kernel.setArg(1, outputImage);
        kernel.setArg(2, width);
        kernel.setArg(3, height);

        // 커널 실행
        cl::NDRange globalSize(width, height);
        queue.enqueueNDRangeKernel(kernel, cl::NullRange, globalSize);
        queue.finish();

        // 결과 읽기
        std::vector<unsigned char> outputPixels(width*height*4);
        cl::size_t<3> origin;
        cl::size_t<3> region;
        origin[0] = origin[1] = origin[2] = 0;
        region[0] = width; region[1] = height; region[2] = 1;

        queue.enqueueReadImage(outputImage, CL_TRUE, origin, region, 0, 0, outputPixels.data());

        // 결과 저장
        stbi_write_png("output.png", width, height, 4, outputPixels.data(), width*4);

        std::cout << "그레이스케일 변환 완료!\n";
    } catch (cl::Error &e) {
        std::cerr << "OpenCL 에러: " << e.what() << " (" << e.err() << ")\n";
        return 1;
    }

    return 0;
}

위 코드 실행 후 output.png를 열어보면 원본 이미지가 그레이스케일로 변환되어 있을 것입니다. CUDA 텍스처 메모리를 다뤄본 경험이 있다면, 전반적인 흐름(이미지를 디바이스에 올리고, 커널에서 처리하고, 결과를 내려받는 과정)이 매우 비슷하다는 것을 느끼실 거예요.

5. CUDA와 비교

• CUDA: 텍스처 메모리(cudaCreateTextureObject, cudaCreateSurfaceObject)를 별도로 설정하고, 텍스처 접근 함수를 통해 픽셀 읽기/쓰기를 수행
• OpenCL: cl::Image2D 객체와 read_imagef, write_imagef 함수 사용. 런타임에 하드웨어별 최적화가 적용될 수 있으며, 다양한 디바이스에서 이미지 처리를 동일한 코드로 수행 가능

결국 개념은 유사하지만, OpenCL은 다양한 디바이스 대응을 위해 조금 더 일반화된 인터페이스를 제공합니다.

6. 참고할만한 자료

• Khronos OpenCL Reference Pages (Images) : OpenCL 이미지 관련 API 문서
• stb 라이브러리 : 이미지 로딩 및 저장을 위한 간단한 헤더-only 라이브러리

7. 마무리

이번 글에서는 OpenCL 이미지 객체를 이용해 실제로 유용한 이미지 처리 예제를 구현해봤습니다. 이제는 단순한 벡터 연산을 넘어, 2D 이미지 처리를 통해 GPU 가속의 장점을 실감하실 수 있을 겁니다.

다음 글에서는 성능 최적화 기초를 다루며, 워크그룹 설정, 메모리 패턴 개선, 프로파일링 등을 가볍게 훑어보겠습니다.

 

유용한 링크 & 리소스

• OpenCL 공식 문서 : OpenCL 표준 및 레퍼런스
• stb_image : 이미지 로딩/저장 라이브러리
• C++ Reference : 최신 C++ 문법 참고