[Vulkan으로 GPGPU 시작하기] #6: Compute 셰이더, 파이프라인 구성, 디스크립터 사용법

지난 글(#5)에서 우리는 Vulkan에서 GPU 메모리를 관리하고, 버퍼나 이미지를 생성하고, 이를 GPU 메모리에 할당하는 기초를 다뤘습니다. 이제는 GPGPU의 핵심인 Compute 셰이더(Compute Shader)를 사용하기 위한 준비 단계로 넘어가겠습니다. Vulkan에서 Compute 셰이더를 실행하려면, 셰이더 코드를 SPIR-V로 컴파일하고, 이를 Compute 파이프라인(Compute Pipeline)에 등록한 뒤, 디스크립터(Descriptor)를 통해 버퍼, 이미지 등 자원을 셰이더에 연결해야 합니다.

CUDA에서는 .cu 파일에 커널 코드를 작성하고 nvcc로 컴파일한 뒤 <<< >>> 표기법으로 커널을 호출하는 방식에 익숙할 것입니다. 하지만 Vulkan에서는 셰이더를 GLSL/HLSL로 작성한 뒤 SPIR-V 중간 표현으로 컴파일하고, Compute 파이프라인 및 디스크립터 세트를 구성한 뒤 명령 버퍼에 해당 파이프라인과 디스크립터를 바인딩하고 Dispatch 명령을 기록하는 식으로 진행합니다. 꽤 번거로워 보이지만, 이 과정을 이해하면 Vulkan의 유연하고 범용적인 구조를 체감할 수 있습니다.

SPIR-V 셰이더 컴파일 개념

Vulkan에서는 셰이더를 SPIR-V라는 중간 표현 언어로 다룹니다. GLSL 또는 HLSL로 셰이더를 작성한 뒤 glslc(LunarG SDK 포함)나 dxc 등을 사용해 SPIR-V 바이너리로 컴파일할 수 있습니다.

예를 들어, 간단한 Compute 셰이더(shader.comp)가 있다고 가정해봅시다. GLSL 형식으로 작성한다면:

#version 450
layout(local_size_x = 256) in; // 워크그룹 당 256 스레드

layout(std430, binding = 0) buffer DataBuffer {
    float data[];
};

void main() {
    uint idx = gl_GlobalInvocationID.x;
    data[idx] = data[idx] * 2.0; // 단순히 값을 두 배로
}

이 코드는 data라는 float 배열을 받아서 각 요소를 두 배로 만드는 Compute 셰이더입니다. CUDA로 치면 __global__ 커널 함수 안에서 data[idx] = data[idx] * 2.0f; 하는 것과 비슷한 로직입니다.

이 셰이더를 SPIR-V로 컴파일하려면 다음과 같이 할 수 있습니다:

glslc shader.comp -o shader.comp.spv

shader.comp.spv 파일이 SPIR-V 바이너리이며, Vulkan에서 이 파일을 로딩해서 파이프라인을 구성할 수 있습니다.

Compute 파이프라인 생성

Compute 파이프라인은 Compute 셰이더 하나를 기반으로 합니다. 파이프라인을 만들 때는 셰이더 모듈(Shader Module)을 먼저 생성한 뒤, 이를 Compute 파이프라인에 연결합니다.

// 셰이더 바이너리를 읽어오는 가정된 함수
// readFile("shader.comp.spv")로 SPIR-V 바이너리 로딩
std::vector<char> code = readFile("shader.comp.spv");

VkShaderModuleCreateInfo moduleCreateInfo = {};
moduleCreateInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;
moduleCreateInfo.codeSize = code.size();
moduleCreateInfo.pCode = (uint32_t*)code.data();

VkShaderModule computeShaderModule;
vkCreateShaderModule(device, &moduleCreateInfo, NULL, &computeShaderModule);

// Compute 파이프라인 생성
VkPipelineShaderStageCreateInfo stageInfo = {};
stageInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
stageInfo.stage = VK_SHADER_STAGE_COMPUTE_BIT;
stageInfo.module = computeShaderModule;
stageInfo.pName = "main"; // 셰이더 엔트리 포인트 이름

VkComputePipelineCreateInfo pipelineInfo = {};
pipelineInfo.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
pipelineInfo.stage = stageInfo;
// 레이아웃은 디스크립터 레이아웃 등 파이프라인 레이아웃 정보 필요, 아래에서 설명

VkPipelineLayoutCreateInfo layoutInfo = {};
layoutInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;
// 디스크립터 레이아웃을 설정할 것이며, 아직 준비 전이므로 뒤에서 다시 다룰 예정

VkPipelineLayout pipelineLayout;
vkCreatePipelineLayout(device, &layoutInfo, NULL, &pipelineLayout);

pipelineInfo.layout = pipelineLayout;

VkPipeline computePipeline;
if (vkCreateComputePipelines(device, VK_NULL_HANDLE, 1, &pipelineInfo, NULL, &computePipeline) != VK_SUCCESS) {
    fprintf(stderr, "Failed to create compute pipeline!\n");
}

여기서 중요한 포인트는 파이프라인을 만들기 전에 “파이프라인 레이아웃(Pipeline Layout)”을 정의해야 한다는 것입니다. 파이프라인 레이아웃은 디스크립터 세트(Descriptor Set) 레이아웃과 Push Constants 등 파이프라인 전역 상태를 정의하는 객체입니다.

디스크립터(Descriptor)와 디스크립터 세트 레이아웃

셰이더에서 데이터를 참조하려면, 버퍼나 이미지를 셰이더와 연결하는 “디스크립터”를 사용해야 합니다. 디스크립터는 셰이더가 접근할 자원을 정의하며, 디스크립터 세트(Descriptor Set)는 이러한 디스크립터들의 집합입니다.

위의 GLSL 예제에서 layout(std430, binding = 0) buffer DataBuffer 구문은 “binding = 0 위치에 있는 스토리지 버퍼”를 의미합니다. 이를 Vulkan에서 설정하려면 디스크립터 셋 레이아웃을 다음과 같이 정의합니다.

VkDescriptorSetLayoutBinding binding = {};
binding.binding = 0;
binding.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
binding.descriptorCount = 1;
binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;

VkDescriptorSetLayoutCreateInfo setLayoutInfo = {};
setLayoutInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;
setLayoutInfo.bindingCount = 1;
setLayoutInfo.pBindings = &binding;

VkDescriptorSetLayout descriptorSetLayout;
vkCreateDescriptorSetLayout(device, &setLayoutInfo, NULL, &descriptorSetLayout);

이제 파이프라인 레이아웃을 만들 때 이 디스크립터 셋 레이아웃을 포함시키면, 셰이더에서 binding=0에 해당하는 디스크립터로 스토리지 버퍼를 접근할 수 있게 됩니다.

VkPipelineLayoutCreateInfo layoutInfo = {};
layoutInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;
layoutInfo.setLayoutCount = 1;
layoutInfo.pSetLayouts = &descriptorSetLayout;

VkPipelineLayout pipelineLayout;
vkCreatePipelineLayout(device, &layoutInfo, NULL, &pipelineLayout);

파이프라인 생성 시 pipelineInfo.layout = pipelineLayout;로 지정하면, 이 파이프라인은 해당 디스크립터 레이아웃을 사용하게 됩니다.

디스크립터 풀(Descriptor Pool)과 디스크립터 셋(Descriptor Set) 할당

디스크립터 셋을 실제로 할당하려면 디스크립터 풀을 먼저 만들어야 합니다. 디스크립터 풀은 디스크립터 셋 할당에 사용되는 리소스 풀입니다.

VkDescriptorPoolSize poolSize = {};
poolSize.type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
poolSize.descriptorCount = 1;

VkDescriptorPoolCreateInfo poolInfo = {};
poolInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO;
poolInfo.poolSizeCount = 1;
poolInfo.pPoolSizes = &poolSize;
poolInfo.maxSets = 1;

VkDescriptorPool descriptorPool;
vkCreateDescriptorPool(device, &poolInfo, NULL, &descriptorPool);

// 디스크립터 셋 할당
VkDescriptorSetAllocateInfo allocInfo = {};
allocInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO;
allocInfo.descriptorPool = descriptorPool;
allocInfo.descriptorSetCount = 1;
allocInfo.pSetLayouts = &descriptorSetLayout;

VkDescriptorSet descriptorSet;
vkAllocateDescriptorSets(device, &allocInfo, &descriptorSet);

이제 descriptorSet이라는 핸들이 생겼고, 여기에 실제 버퍼를 바인딩할 수 있습니다.

디스크립터 업데이트(Descriptor Update)

버퍼를 디스크립터에 연결하기 위해 vkUpdateDescriptorSets()를 사용합니다. 앞서 만들었던 buffer와 bufferMemory를 바인딩한 버퍼를 디스크립터에 업데이트합니다.

VkDescriptorBufferInfo bufferInfo = {};
bufferInfo.buffer = buffer;      // 이전 글에서 만든 스토리지 버퍼
bufferInfo.offset = 0;
bufferInfo.range = VK_WHOLE_SIZE;

VkWriteDescriptorSet write = {};
write.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;
write.dstSet = descriptorSet;
write.dstBinding = 0; // binding=0과 매치
write.dstArrayElement = 0;
write.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
write.descriptorCount = 1;
write.pBufferInfo = &bufferInfo;

vkUpdateDescriptorSets(device, 1, &write, 0, NULL);

이제 이 descriptorSet은 셰이더가 참조할 버퍼를 알게 되었습니다.

정리

• GLSL/HLSL로 Compute 셰이더 작성 후 SPIR-V로 컴파일
• Shader Module 생성, Compute 파이프라인과 파이프라인 레이아웃 구성
• 디스크립터 셋 레이아웃 정의로 셰이더 바인딩 구조 명시
• 디스크립터 풀에서 디스크립터 셋 할당 및 버퍼 연결

CUDA에서는 .cu 커널을 NVCC로 빌드하고, 실행 시 cudaSetDevice, cudaMemcpy, kernel<<<>>> 호출 등으로 비교적 단순히 처리할 수 있습니다. 반면 Vulkan은 더 많은 설정 과정이 필요하지만, 그래픽 파이프라인과도 유사한 구조로 통합 처리할 수 있고, 다양한 자원 관리 전략을 세울 수 있습니다.

다음 글 예고

다음 글(#7)에서는 지금까지 준비한 파이프라인, 디스크립터, 버퍼를 실제로 명령 버퍼에 기록하고, Compute Shader를 디스패치(Dispatch)하여 GPGPU 연산을 수행하는 과정을 다룰 예정입니다. 이로써 “데이터를 GPU에 올리고, 셰이더를 통해 연산하고, 결과를 확인하는” 전 과정을 완성할 수 있게 됩니다.

유용한 링크 & 리소스

• Vulkan Tutorial - Compute pipeline
• Khronos Vulkan Docs
• CUDA 공식 문서 - CUDA 커널 호출과 비교해 Vulkan 파이프라인 구성 이해