[Vulkan으로 GPGPU 시작하기] #7: 벡터 덧셈 예제로 전체 흐름 익히기

지난 글(#6)에서는 Compute 셰이더를 SPIR-V로 컴파일하고, 이를 기반으로 Compute 파이프라인을 만들고, 디스크립터를 활용해 셰이더에 데이터를 연결하는 방법을 살펴봤습니다. 이제 이 모든 준비 과정을 종합해 실제로 GPGPU 연산을 수행하는 간단한 예제를 만들어보겠습니다. 이번 글에서는 벡터 덧셈(Vector Addition) 예제를 통해, 호스트에서 데이터를 준비하고 GPU로 전달한 뒤, Compute 셰이더를 통해 연산을 수행하고 결과를 다시 가져오는 전체 흐름을 정리해볼 것입니다.

CUDA에서 벡터 덧셈은 대략 cudaMalloc, cudaMemcpy, kernel<<<>>>, cudaMemcpy라는 단순한 과정을 거치며, 커널 코드도 data[idx] = a[idx] + b[idx];처럼 직관적입니다. 반면 Vulkan에서는 인스턴스 생성부터 디바이스, 큐, 커맨드 버퍼, 메모리 할당, 파이프라인, 디스크립터 셋 준비 등을 모두 거친 뒤에야 비로소 vkQueueSubmit()로 연산을 실행합니다. 이번 예제는 이 복잡한 과정을 작은 코드 예제로 묶어, 실제로 GPGPU 연산이 어떻게 Vulkan에서 이루어지는지 감을 잡을 수 있게 해줄 것입니다.

예제 개요

우리가 구현할 벡터 덧셈 과정은 다음과 같습니다.

1. 호스트에서 두 개의 float 배열 A, B를 준비하고, GPU에 전달할 버퍼를 만듬.
2. A와 B를 GPU 메모리에 올려놓고, 결과를 저장할 버퍼 C도 준비.
3. Compute 셰이더를 통해 C[idx] = A[idx] + B[idx] 연산 수행.
4. 연산 결과를 다시 호스트로 가져와서 C 값을 출력해 확인.

이 예제를 통해 Vulkan의 초기화부터 커맨드 버퍼, 파이프라인, 디스크립터, 메모리 관리, 연산 실행까지 한 번에 체득할 수 있을 것입니다.

셰이더 코드 (GLSL)

간단한 Compute 셰이더를 준비합니다. add.comp라는 GLSL 셰이더 파일을 예로 들면:

#version 450
layout(local_size_x = 256) in;

layout(std430, binding = 0) buffer ABuffer {
    float A[];
};

layout(std430, binding = 1) buffer BBuffer {
    float B[];
};

layout(std430, binding = 2) buffer CBuffer {
    float C[];
};

void main() {
    uint idx = gl_GlobalInvocationID.x;
    C[idx] = A[idx] + B[idx];
}

이 셰이더는 A, B, C 버퍼를 입력받아 A+B를 C에 저장하는 단순한 연산을 합니다. gl_GlobalInvocationID.x는 스레드 인덱스를 의미합니다. CUDA로 치면 blockIdx.x * blockDim.x + threadIdx.x에 해당하는 인덱싱과 유사한 개념입니다.

이 셰이더를 컴파일해 SPIR-V를 얻습니다.

glslc add.comp -o add.comp.spv

호스트 코드 개요

코드 전체를 여기서 모두 나열하기는 길겠지만, 주요 흐름을 정리해보겠습니다.

1. Vulkan 초기화: 인스턴스, 물리 디바이스 선택, 로지컬 디바이스, 큐 준비 (이전 글들 참고)
2. 버퍼 생성 및 메모리 할당:
   • float 배열 A, B를 준비하고, C를 담을 버퍼 준비
   • A, B, C용 버퍼를 각각 VK_BUFFER_USAGE_STORAGE_BUFFER_BIT로 만들고 메모리 할당
   • A와 B 데이터는 Host Visible 메모리에 올려 GPU로 복사 (매핑 후 memcpy)
3. 디스크립터 및 파이프라인 설정:
   • binding=0,1,2에 각각 A, B, C 버퍼를 바인딩할 디스크립터 셋 레이아웃 정의
   • Compute 파이프라인 생성 (add.comp.spv 셰이더 사용)
   • 디스크립터 풀, 디스크립터 셋 할당 후 vkUpdateDescriptorSets()로 A,B,C 버퍼 연결
4. 커맨드 버퍼에 명령 기록:
   • vkCmdBindPipeline()로 Compute 파이프라인 바인딩
   • vkCmdBindDescriptorSets()로 디스크립터 셋 바인딩
   • vkCmdDispatch()로 스레드 실행 (예: 벡터 길이가 N이라면 dispatch((N+255)/256,1,1) 형태)
5. vkQueueSubmit()와 vkQueueWaitIdle()로 연산 실행
6. 결과 회수: C 버퍼를 다시 호스트 메모리에 매핑해 C 값 출력

이 모든 과정을 순차적으로 수행하면, GPU를 통해 A+B 결과를 C에 담아 확인할 수 있습니다.

실제 코드 흐름 예제 (간략)

아래는 핵심 단계만 추린 의사 코드입니다. (에러 처리, 리소스 해제 생략)

// 1. Vulkan 초기화 (인스턴스, 디바이스, 큐) - 이전 글 예제 코드 활용
// 2. A, B, C 데이터 준비
int N = 1024;
std::vector<float> A(N), B(N), C(N, 0);
for (int i = 0; i < N; i++) {
    A[i] = (float)i;
    B[i] = (float)(N - i);
}

// 3. 버퍼 생성 및 메모리 할당 (A_buffer, B_buffer, C_buffer)
// Host Visible 메모리에 할당하고 A,B 값을 memcpy로 올림

// 4. 디스크립터 셋 레이아웃, 파이프라인 레이아웃, Compute 파이프라인 생성
// 셰이더 모듈 생성(add.comp.spv), descriptorSetLayout, pipelineLayout, computePipeline 준비

// 5. 디스크립터 풀, 디스크립터 셋 할당 및 업데이트
// binding=0 -> A_buffer, binding=1 -> B_buffer, binding=2 -> C_buffer 설정

// 6. 커맨드 버퍼 레코딩
vkBeginCommandBuffer(commandBuffer, &beginInfo);
vkCmdBindPipeline(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, computePipeline);
vkCmdBindDescriptorSets(commandBuffer, VK_PIPELINE_BIND_POINT_COMPUTE,
    pipelineLayout, 0, 1, &descriptorSet, 0, NULL);

// 스레드 개수 N, local_size_x=256 이므로 dispatch((N+255)/256,1,1)
uint32_t groupCountX = (N + 255) / 256;
vkCmdDispatch(commandBuffer, groupCountX, 1, 1);

vkEndCommandBuffer(commandBuffer);

// 7. 큐 제출
vkQueueSubmit(computeQueue, 1, &submitInfo, VK_NULL_HANDLE);
vkQueueWaitIdle(computeQueue);

// 8. 결과 회수
// C_buffer를 다시 매핑해 C 배열을 가져와 값 검사
vkMapMemory(device, C_bufferMemory, 0, VK_WHOLE_SIZE, 0, &mappedPtr);
float* mappedData = (float*)mappedPtr;
for (int i = 0; i < 5; i++) {
    printf("C[%d] = %f\n", i, mappedData[i]); // 기대값: A[i]+B[i] = i+(N-i) = N
}
vkUnmapMemory(device, C_bufferMemory);

이렇게 하면 C 벡터는 모두 N이라는 값이 되어야 합니다. 예를 들어 N=1024라면 C[i]=1024.0f가 나와야 합니다.

CUDA와 비교하며 느끼는 점

CUDA에서는 이 과정이 훨씬 짧습니다. 디바이스 선택 후 cudaMalloc, cudaMemcpy로 데이터 전송, kernel<<<>>> 호출로 연산, 다시 cudaMemcpy로 결과 회수가 끝입니다. Vulkan은 훨씬 더 복잡한 초기화와 설정 과정이 필요하지만, 그만큼 그래픽 파이프라인 및 다양한 자원 관리 패턴과 융합이 가능하고, 하드웨어 독립적이며, 멀티 디바이스나 고급 최적화에도 유연하게 대처할 수 있습니다.

이 벡터 덧셈 예제는 Vulkan GPGPU의 “Hello World!” 같은 것이며, 이 과정을 통해 Vulkan의 GPGPU 워크플로우 전반을 한 번에 체득할 수 있습니다.

다음 글 예고

다음 글(#8)에서는 이제 디버깅, Validation Layer, 성능 프로파일링 같은 추가적인 개발 지원 툴들에 대해 간단히 살펴보겠습니다. 복잡한 Vulkan 환경에서 문제가 발생했을 때 이를 어떻게 추적하고 해결할지, 성능 최적화를 위해 어떤 도구를 사용할 수 있을지 개략적으로 소개할 예정입니다.

유용한 링크 & 리소스

• Vulkan Tutorial - Compute Example
• NVIDIA Developer Blogs - Vulkan Compute
• CUDA 공식 문서 - CUDA 벡터 덧셈 예제와 비교하며 이해