Vulkan Compute và GPU đa nền tảng

Vulkan là API tính toán GPU duy nhất chạy trên mọi nền tảng lớn: NVIDIA, AMD, Intel, Apple (qua MoltenVK), Android, và thậm chí cả trình duyệt (qua WebGPU). File này đề cập đến kiến trúc Vulkan, pipeline tính toán, viết compute shader trong GLSL, thiết lập C++ đầy đủ cho một chương trình tính toán GPU, bộ nhớ dùng chung và đồng bộ hóa, WebGPU cho trình duyệt, và các ví dụ suy luận ML thực tế.

• CUDA thống trị huấn luyện ML trên phần cứng NVIDIA. Nhưng không phải mục tiêu triển khai nào cũng có GPU NVIDIA. Một ứng dụng mobile chạy trên GPU Qualcomm Adreno hoặc ARM Mali. Một ứng dụng web chạy trong trình duyệt. Một engine game cần hỗ trợ đồng thời AMD, Intel và NVIDIA. Cho tất cả những trường hợp này, Vulkan là câu trả lời.

• Vulkan rất dài dòng — một chương trình tính toán "hello world" tốn ~300 dòng C++. Nhưng sự dài dòng này là cái giá của sự kiểm soát tường minh: bạn tự quản lý mọi tài nguyên GPU (bộ nhớ, pipeline, command buffer). Sự kiểm soát này cho phép hiệu năng và tính di động tối đa, đánh đổi bằng tốc độ phát triển.

Tổng quan kiến trúc Vulkan

• Vulkan là một API GPU mức thấp do Khronos Group tạo ra (tổ chức đứng sau OpenGL). Khác với CUDA (che giấu việc quản lý tài nguyên GPU), Vulkan yêu cầu bạn quản lý tường minh:

  • Instance và device: tạo một Vulkan instance, liệt kê các GPU khả dụng, và chọn một cái.
  • Bộ nhớ: cấp phát bộ nhớ GPU tường minh, chỉ định loại bộ nhớ (device-local để nhanh, host-visible để CPU truy cập).
  • Buffer: tạo các đối tượng buffer tham chiếu đến bộ nhớ đã cấp phát.
  • Descriptor sets: gắn các buffer vào đầu vào của shader (giống như tham số hàm cho compute shader).
  • Compute pipeline: biên dịch shader và tạo một đối tượng pipeline.
  • Command buffer: ghi lại một chuỗi các lệnh GPU (gắn pipeline, gắn descriptor, dispatch tính toán).
  • Queue submission: nộp command buffer cho GPU để thực thi.
  • Đồng bộ hóa: fence và barrier để đảm bảo thứ tự thực thi đúng.

• Đây là một sự khác biệt căn bản so với mô hình cudaMalloc + khởi chạy kernel của CUDA. Trong CUDA, driver xử lý hầu hết việc này đằng sau hậu trường. Trong Vulkan, bạn tự làm lấy.

Tại sao lại dài dòng?

• Sự tường minh của Vulkan tồn tại vì hai lý do:

  1. Đơn giản hóa driver: các driver OpenGL cực kỳ phức tạp (chúng phải đoán ý định của ứng dụng và tối ưu hóa tương ứng). Vulkan chuyển trách nhiệm đó sang ứng dụng, làm cho driver mỏng hơn, dễ dự đoán hơn, và dễ cài đặt đúng trên nhiều hãng hơn.

  2. Hiệu năng: kiểm soát tường minh bố cục bộ nhớ, đồng bộ hóa, và gộp lệnh cho phép ứng dụng đưa ra các quyết định tối ưu. Trong CUDA, driver có thể chèn các đồng bộ hóa không cần thiết. Trong Vulkan, bạn chỉ đồng bộ khi cần.

Compute Shader trong GLSL

• Compute shader là một chương trình chạy trên GPU, tương tự một kernel CUDA. Nó được viết bằng GLSL (OpenGL Shading Language) và biên dịch thành bytecode SPIR-V (một định dạng nhị phân di động).

Phép cộng Vector

// add.comp — biên dịch với: glslangValidator -V add.comp -o add.spv
#version 450

// Kích thước workgroup: 256 invocation mỗi workgroup (= số luồng mỗi khối trong CUDA)
layout(local_size_x = 256) in;

// Buffer bindings (giống như tham số kernel)
layout(set = 0, binding = 0) buffer InputA { float a[]; };
layout(set = 0, binding = 1) buffer InputB { float b[]; };
layout(set = 0, binding = 2) buffer Output { float c[]; };

// Push constant: dữ liệu uniform nhỏ (giống như tham số kernel)
layout(push_constant) uniform PushConstants {
    uint n;  // số lượng phần tử
};

void main() {
    uint idx = gl_GlobalInvocationID.x;  // chỉ số luồng toàn cục
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

• Ánh xạ sang các khái niệm CUDA:

Vulkan	CUDA	Ý nghĩa
Workgroup	Block	Nhóm các luồng có thể chia sẻ bộ nhớ
Invocation	Thread	Đơn vị thực thi đơn
gl_GlobalInvocationID	blockIdx * blockDim + threadIdx	Chỉ số luồng toàn cục
gl_LocalInvocationID	threadIdx	Chỉ số luồng trong workgroup
gl_WorkGroupID	blockIdx	Chỉ số workgroup
local_size_x	blockDim.x	Số luồng mỗi workgroup
Storage buffer	Global memory	Bộ nhớ GPU đọc/ghi
Shared memory (shared)	__shared__	Bộ nhớ nhanh mỗi workgroup
Push constant	Kernel argument	Dữ liệu uniform nhỏ

ReLU với Bộ nhớ dùng chung

// relu_shared.comp
#version 450

layout(local_size_x = 256) in;

layout(set = 0, binding = 0) buffer Input  { float input_data[]; };
layout(set = 0, binding = 1) buffer Output { float output_data[]; };

layout(push_constant) uniform PushConstants { uint n; };

// Bộ nhớ dùng chung (tương đương CUDA __shared__)
shared float tile[256];

void main() {
    uint gid = gl_GlobalInvocationID.x;
    uint lid = gl_LocalInvocationID.x;

    // Nạp vào bộ nhớ dùng chung
    if (gid < n) {
        tile[lid] = input_data[gid];
    }

    // Barrier: chờ tất cả invocation trong workgroup nạp xong
    barrier();  // tương đương CUDA __syncthreads()

    // Tính ReLU
    if (gid < n) {
        output_data[gid] = max(tile[lid], 0.0);
    }
}

• Với ReLU, bộ nhớ dùng chung không thực sự cần thiết (phép toán theo từng phần tử). Nhưng đây minh họa mẫu: nạp vào bộ nhớ dùng chung → barrier → tính toán → lưu. Với các phép toán cần dữ liệu từ các luồng lân cận (tích chập, reduction, softmax), bộ nhớ dùng chung là thiết yếu.

Reduction Song song (Tổng)

// reduce_sum.comp
#version 450

layout(local_size_x = 256) in;

layout(set = 0, binding = 0) buffer Input  { float input_data[]; };
layout(set = 0, binding = 1) buffer Output { float partial_sums[]; };

layout(push_constant) uniform PushConstants { uint n; };

shared float sdata[256];

void main() {
    uint gid = gl_GlobalInvocationID.x;
    uint lid = gl_LocalInvocationID.x;
    uint wgid = gl_WorkGroupID.x;

    // Nạp vào bộ nhớ dùng chung
    sdata[lid] = (gid < n) ? input_data[gid] : 0.0;
    barrier();

    // Tree reduction trong workgroup
    for (uint stride = 128; stride > 0; stride >>= 1) {
        if (lid < stride) {
            sdata[lid] += sdata[lid + stride];
        }
        barrier();
    }

    // Luồng 0 ghi tổng cục bộ của workgroup
    if (lid == 0) {
        partial_sums[wgid] = sdata[0];
    }
}

• Đây là mẫu reduction song song kinh điển (giống CUDA). Mỗi workgroup tạo ra một tổng cục bộ. Một lần dispatch thứ hai gộp các tổng cục bộ thành kết quả cuối. Tree reduction giảm một nửa số luồng hoạt động ở mỗi bước: 256 → 128 → 64 → ... → 1.

Nhân ma trận với Tiling

// matmul_tiled.comp
#version 450

#define TILE_SIZE 16

layout(local_size_x = TILE_SIZE, local_size_y = TILE_SIZE) in;

layout(set = 0, binding = 0) buffer MatA { float A[]; };
layout(set = 0, binding = 1) buffer MatB { float B[]; };
layout(set = 0, binding = 2) buffer MatC { float C[]; };

layout(push_constant) uniform PushConstants {
    uint M, N, K;
};

shared float tileA[TILE_SIZE][TILE_SIZE];
shared float tileB[TILE_SIZE][TILE_SIZE];

void main() {
    uint row = gl_GlobalInvocationID.y;
    uint col = gl_GlobalInvocationID.x;
    uint lr = gl_LocalInvocationID.y;
    uint lc = gl_LocalInvocationID.x;

    float sum = 0.0;

    for (uint t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // Nạp tile của A và B vào bộ nhớ dùng chung
        uint aCol = t * TILE_SIZE + lc;
        uint bRow = t * TILE_SIZE + lr;

        tileA[lr][lc] = (row < M && aCol < K) ? A[row * K + aCol] : 0.0;
        tileB[lr][lc] = (bRow < K && col < N) ? B[bRow * N + col] : 0.0;

        barrier();

        // Tính tích vô hướng cục bộ
        for (uint k = 0; k < TILE_SIZE; k++) {
            sum += tileA[lr][k] * tileB[k][lc];
        }

        barrier();
    }

    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

• Đây là cùng thuật toán tiling với bản CUDA (file 04), chỉ khác cú pháp GLSL. Các khái niệm hoàn toàn giống nhau: nạp tile vào bộ nhớ dùng chung, barrier, tính toán, barrier, lặp lại.

Thiết lập Vulkan bằng C++

• Compute shader là phần dễ. Phần khó là đoạn mã nền (boilerplate) C++ tạo instance Vulkan, cấp phát bộ nhớ, gắn buffer, và nộp lệnh. Dưới đây là bản tóm gọn của toàn bộ pipeline:

// vulkan_compute.cpp — một ví dụ Vulkan compute tối thiểu nhưng hoàn chỉnh
// Biên dịch: g++ -O3 -o vulkan_compute vulkan_compute.cpp -lvulkan
// Yêu cầu: Vulkan SDK đã cài, add.spv biên dịch từ add.comp

#include <vulkan/vulkan.h>
#include <iostream>
#include <vector>
#include <fstream>
#include <cassert>

// Hàm phụ: đọc file SPIR-V
std::vector<uint32_t> readSPIRV(const std::string& filename) {
    std::ifstream file(filename, std::ios::ate | std::ios::binary);
    size_t fileSize = file.tellg();
    std::vector<uint32_t> buffer(fileSize / sizeof(uint32_t));
    file.seekg(0);
    file.read(reinterpret_cast<char*>(buffer.data()), fileSize);
    return buffer;
}

int main() {
    const uint32_t N = 1024;
    const size_t bufferSize = N * sizeof(float);

    // ========== 1. Tạo Vulkan Instance ==========
    VkApplicationInfo appInfo{};
    appInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;
    appInfo.apiVersion = VK_API_VERSION_1_2;

    VkInstanceCreateInfo instanceInfo{};
    instanceInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;
    instanceInfo.pApplicationInfo = &appInfo;

    VkInstance instance;
    vkCreateInstance(&instanceInfo, nullptr, &instance);

    // ========== 2. Chọn Physical Device (GPU) ==========
    uint32_t deviceCount = 0;
    vkEnumeratePhysicalDevices(instance, &deviceCount, nullptr);
    std::vector<VkPhysicalDevice> devices(deviceCount);
    vkEnumeratePhysicalDevices(instance, &deviceCount, devices.data());
    VkPhysicalDevice physicalDevice = devices[0];  // dùng GPU đầu tiên

    // In tên GPU
    VkPhysicalDeviceProperties props;
    vkGetPhysicalDeviceProperties(physicalDevice, &props);
    std::cout << "Đang dùng GPU: " << props.deviceName << "\n";

    // ========== 3. Tìm Compute Queue Family ==========
    uint32_t queueFamilyCount = 0;
    vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &queueFamilyCount, nullptr);
    std::vector<VkQueueFamilyProperties> queueFamilies(queueFamilyCount);
    vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &queueFamilyCount, queueFamilies.data());

    uint32_t computeFamily = 0;
    for (uint32_t i = 0; i < queueFamilyCount; i++) {
        if (queueFamilies[i].queueFlags & VK_QUEUE_COMPUTE_BIT) {
            computeFamily = i;
            break;
        }
    }

    // ========== 4. Tạo Logical Device và Queue ==========
    float queuePriority = 1.0f;
    VkDeviceQueueCreateInfo queueInfo{};
    queueInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;
    queueInfo.queueFamilyIndex = computeFamily;
    queueInfo.queueCount = 1;
    queueInfo.pQueuePriorities = &queuePriority;

    VkDeviceCreateInfo deviceInfo{};
    deviceInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;
    deviceInfo.queueCreateInfoCount = 1;
    deviceInfo.pQueueCreateInfos = &queueInfo;

    VkDevice device;
    vkCreateDevice(physicalDevice, &deviceInfo, nullptr, &device);

    VkQueue computeQueue;
    vkGetDeviceQueue(device, computeFamily, 0, &computeQueue);

    // ========== 5. Cấp phát Buffer (A, B, C) ==========
    // Để ngắn gọn, dùng host-visible memory (chậm hơn nhưng đơn giản hơn)
    auto createBuffer = [&](VkBuffer& buffer, VkDeviceMemory& memory) {
        VkBufferCreateInfo bufInfo{};
        bufInfo.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
        bufInfo.size = bufferSize;
        bufInfo.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT;
        vkCreateBuffer(device, &bufInfo, nullptr, &buffer);

        VkMemoryRequirements memReqs;
        vkGetBufferMemoryRequirements(device, buffer, &memReqs);

        // Tìm loại bộ nhớ host-visible
        VkPhysicalDeviceMemoryProperties memProps;
        vkGetPhysicalDeviceMemoryProperties(physicalDevice, &memProps);
        uint32_t memType = 0;
        for (uint32_t i = 0; i < memProps.memoryTypeCount; i++) {
            if ((memReqs.memoryTypeBits & (1 << i)) &&
                (memProps.memoryTypes[i].propertyFlags &
                 (VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT))) {
                memType = i;
                break;
            }
        }

        VkMemoryAllocateInfo allocInfo{};
        allocInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
        allocInfo.allocationSize = memReqs.size;
        allocInfo.memoryTypeIndex = memType;
        vkAllocateMemory(device, &allocInfo, nullptr, &memory);
        vkBindBufferMemory(device, buffer, memory, 0);
    };

    VkBuffer bufA, bufB, bufC;
    VkDeviceMemory memA, memB, memC;
    createBuffer(bufA, memA);
    createBuffer(bufB, memB);
    createBuffer(bufC, memC);

    // ========== 6. Điền Buffer đầu vào ==========
    float* ptrA;
    vkMapMemory(device, memA, 0, bufferSize, 0, (void**)&ptrA);
    for (uint32_t i = 0; i < N; i++) ptrA[i] = 1.0f;
    vkUnmapMemory(device, memA);

    float* ptrB;
    vkMapMemory(device, memB, 0, bufferSize, 0, (void**)&ptrB);
    for (uint32_t i = 0; i < N; i++) ptrB[i] = 2.0f;
    vkUnmapMemory(device, memB);

    // ========== 7. Tạo Compute Pipeline ==========
    auto spirvCode = readSPIRV("add.spv");
    VkShaderModuleCreateInfo shaderInfo{};
    shaderInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;
    shaderInfo.codeSize = spirvCode.size() * sizeof(uint32_t);
    shaderInfo.pCode = spirvCode.data();
    VkShaderModule shaderModule;
    vkCreateShaderModule(device, &shaderInfo, nullptr, &shaderModule);

    // Descriptor set layout (cho Vulkan biết về các buffer binding)
    VkDescriptorSetLayoutBinding bindings[3] = {};
    for (int i = 0; i < 3; i++) {
        bindings[i].binding = i;
        bindings[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
        bindings[i].descriptorCount = 1;
        bindings[i].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
    }

    VkDescriptorSetLayoutCreateInfo layoutInfo{};
    layoutInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;
    layoutInfo.bindingCount = 3;
    layoutInfo.pBindings = bindings;
    VkDescriptorSetLayout descLayout;
    vkCreateDescriptorSetLayout(device, &layoutInfo, nullptr, &descLayout);

    // Push constant range
    VkPushConstantRange pushRange{};
    pushRange.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
    pushRange.offset = 0;
    pushRange.size = sizeof(uint32_t);

    // Pipeline layout
    VkPipelineLayoutCreateInfo pipeLayoutInfo{};
    pipeLayoutInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;
    pipeLayoutInfo.setLayoutCount = 1;
    pipeLayoutInfo.pSetLayouts = &descLayout;
    pipeLayoutInfo.pushConstantRangeCount = 1;
    pipeLayoutInfo.pPushConstantRanges = &pushRange;
    VkPipelineLayout pipelineLayout;
    vkCreatePipelineLayout(device, &pipeLayoutInfo, nullptr, &pipelineLayout);

    // Compute pipeline
    VkComputePipelineCreateInfo pipeInfo{};
    pipeInfo.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
    pipeInfo.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
    pipeInfo.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;
    pipeInfo.stage.module = shaderModule;
    pipeInfo.stage.pName = "main";
    pipeInfo.layout = pipelineLayout;
    VkPipeline pipeline;
    vkCreateComputePipelines(device, VK_NULL_HANDLE, 1, &pipeInfo, nullptr, &pipeline);

    // ========== 8. Descriptor Set (gắn buffer vào shader) ==========
    VkDescriptorPoolSize poolSize{};
    poolSize.type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
    poolSize.descriptorCount = 3;

    VkDescriptorPoolCreateInfo poolInfo{};
    poolInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO;
    poolInfo.maxSets = 1;
    poolInfo.poolSizeCount = 1;
    poolInfo.pPoolSizes = &poolSize;
    VkDescriptorPool descPool;
    vkCreateDescriptorPool(device, &poolInfo, nullptr, &descPool);

    VkDescriptorSetAllocateInfo descAllocInfo{};
    descAllocInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO;
    descAllocInfo.descriptorPool = descPool;
    descAllocInfo.descriptorSetCount = 1;
    descAllocInfo.pSetLayouts = &descLayout;
    VkDescriptorSet descSet;
    vkAllocateDescriptorSets(device, &descAllocInfo, &descSet);

    // Ghi tham chiếu buffer vào descriptor set
    VkDescriptorBufferInfo bufInfos[3] = {
        {bufA, 0, bufferSize}, {bufB, 0, bufferSize}, {bufC, 0, bufferSize}
    };
    VkWriteDescriptorSet writes[3] = {};
    for (int i = 0; i < 3; i++) {
        writes[i].sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;
        writes[i].dstSet = descSet;
        writes[i].dstBinding = i;
        writes[i].descriptorCount = 1;
        writes[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
        writes[i].pBufferInfo = &bufInfos[i];
    }
    vkUpdateDescriptorSets(device, 3, writes, 0, nullptr);

    // ========== 9. Ghi và Nộp Command Buffer ==========
    VkCommandPoolCreateInfo cmdPoolInfo{};
    cmdPoolInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
    cmdPoolInfo.queueFamilyIndex = computeFamily;
    VkCommandPool cmdPool;
    vkCreateCommandPool(device, &cmdPoolInfo, nullptr, &cmdPool);

    VkCommandBufferAllocateInfo cmdAllocInfo{};
    cmdAllocInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
    cmdAllocInfo.commandPool = cmdPool;
    cmdAllocInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
    cmdAllocInfo.commandBufferCount = 1;
    VkCommandBuffer cmdBuf;
    vkAllocateCommandBuffers(device, &cmdAllocInfo, &cmdBuf);

    VkCommandBufferBeginInfo beginInfo{};
    beginInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
    vkBeginCommandBuffer(cmdBuf, &beginInfo);

    vkCmdBindPipeline(cmdBuf, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);
    vkCmdBindDescriptorSets(cmdBuf, VK_PIPELINE_BIND_POINT_COMPUTE,
                            pipelineLayout, 0, 1, &descSet, 0, nullptr);
    vkCmdPushConstants(cmdBuf, pipelineLayout, VK_SHADER_STAGE_COMPUTE_BIT,
                       0, sizeof(uint32_t), &N);
    vkCmdDispatch(cmdBuf, (N + 255) / 256, 1, 1);  // khởi chạy workgroup

    vkEndCommandBuffer(cmdBuf);

    // Nộp
    VkFenceCreateInfo fenceInfo{};
    fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO;
    VkFence fence;
    vkCreateFence(device, &fenceInfo, nullptr, &fence);

    VkSubmitInfo submitInfo{};
    submitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
    submitInfo.commandBufferCount = 1;
    submitInfo.pCommandBuffers = &cmdBuf;
    vkQueueSubmit(computeQueue, 1, &submitInfo, fence);
    vkWaitForFences(device, 1, &fence, VK_TRUE, UINT64_MAX);

    // ========== 10. Đọc kết quả ==========
    float* ptrC;
    vkMapMemory(device, memC, 0, bufferSize, 0, (void**)&ptrC);
    std::cout << "Kết quả: c[0]=" << ptrC[0] << " c[1]=" << ptrC[1]
              << " (kỳ vọng 3.0)\n";
    bool correct = true;
    for (uint32_t i = 0; i < N; i++) {
        if (ptrC[i] != 3.0f) { correct = false; break; }
    }
    std::cout << (correct ? "TẤT CẢ ĐÚNG" : "CÓ LỖI") << "\n";
    vkUnmapMemory(device, memC);

    // ========== Dọn dẹp (rút gọn) ==========
    vkDestroyFence(device, fence, nullptr);
    vkDestroyCommandPool(device, cmdPool, nullptr);
    vkDestroyPipeline(device, pipeline, nullptr);
    vkDestroyPipelineLayout(device, pipelineLayout, nullptr);
    vkDestroyDescriptorPool(device, descPool, nullptr);
    vkDestroyDescriptorSetLayout(device, descLayout, nullptr);
    vkDestroyShaderModule(device, shaderModule, nullptr);
    vkDestroyBuffer(device, bufA, nullptr); vkFreeMemory(device, memA, nullptr);
    vkDestroyBuffer(device, bufB, nullptr); vkFreeMemory(device, memB, nullptr);
    vkDestroyBuffer(device, bufC, nullptr); vkFreeMemory(device, memC, nullptr);
    vkDestroyDevice(device, nullptr);
    vkDestroyInstance(instance, nullptr);

    return 0;
}

• Đúng vậy, đây là ~200 dòng chỉ cho phép cộng vector. So với ~30 dòng của CUDA. Đây là cái giá của sự tường minh. Nhưng hãy lưu ý: mọi dòng đều có mục đích. Không có quyết định driver ẩn, không có đồng bộ hóa ngầm, không có cấp phát bất ngờ. Bạn kiểm soát tất cả.

• Trong thực tế, bạn sẽ bọc đoạn mã nền này vào một thư viện hỗ trợ (hoặc dùng một thư viện có sẵn như vk-bootstrap, VMA cho cấp phát bộ nhớ, hoặc kompute cho tính toán Vulkan tập trung vào ML).

Kompute: Vulkan đơn giản hóa cho ML

• Kompute là một thư viện C++ mã nguồn mở bọc lấy đoạn mã nền của Vulkan cho tính toán GPU. Cùng phép cộng vector trở thành:

#include <kompute/Kompute.hpp>

int main() {
    kp::Manager mgr;

    auto tensorA = mgr.tensor({1, 1, 1, 1, 1});
    auto tensorB = mgr.tensor({2, 2, 2, 2, 2});
    auto tensorC = mgr.tensor({0, 0, 0, 0, 0});

    std::string shader = R"(
        #version 450
        layout(local_size_x = 1) in;
        layout(set=0, binding=0) buffer A { float a[]; };
        layout(set=0, binding=1) buffer B { float b[]; };
        layout(set=0, binding=2) buffer C { float c[]; };
        void main() {
            uint i = gl_GlobalInvocationID.x;
            c[i] = a[i] + b[i];
        }
    )";

    auto algorithm = mgr.algorithm({tensorA, tensorB, tensorC},
                                     kompute::Shader::compile_source(shader));

    mgr.sequence()
        ->record<kp::OpTensorSyncDevice>({tensorA, tensorB, tensorC})
        ->record<kp::OpAlgoDispatch>(algorithm)
        ->record<kp::OpTensorSyncLocal>({tensorC})
        ->eval();

    // tensorC bây giờ chứa [3, 3, 3, 3, 3]
}

• Dễ đọc hơn nhiều. Kompute xử lý việc tạo instance, chọn device, cấp phát bộ nhớ, descriptor set, và quản lý command buffer. Bạn tập trung vào shader và dữ liệu.

WebGPU: Tính toán GPU trong trình duyệt

• WebGPU là người kế nhiệm của WebGL, cung cấp khả năng truy cập GPU hiện đại từ JavaScript. Nó được xây dựng trên Vulkan (Linux/Android), Metal (macOS/iOS), và DirectX 12 (Windows), che giấu các khác biệt nền tảng.

• WebGPU dùng WGSL (WebGPU Shading Language) thay vì GLSL:

// add.wgsl — WebGPU compute shader
@group(0) @binding(0) var<storage, read> a: array<f32>;
@group(0) @binding(1) var<storage, read> b: array<f32>;
@group(0) @binding(2) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
    let i = id.x;
    c[i] = a[i] + b[i];
}

• Thiết lập JavaScript (tóm gọn):

const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();

// Tạo buffer
const bufferA = device.createBuffer({ size: N * 4, usage: GPUBufferUsage.STORAGE, mappedAtCreation: true });
new Float32Array(bufferA.getMappedRange()).fill(1.0);
bufferA.unmap();

// ... (tương tự cho B và C)

// Tạo pipeline từ shader WGSL
const pipeline = device.createComputePipeline({
    layout: 'auto',
    compute: { module: device.createShaderModule({ code: wgslSource }), entryPoint: 'main' }
});

// Dispatch
const encoder = device.createCommandEncoder();
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(Math.ceil(N / 256));
pass.end();
device.queue.submit([encoder.finish()]);

• Tại sao WebGPU quan trọng với ML: chạy suy luận trong trình duyệt nghĩa là không tốn chi phí server, không độ trễ, và dữ liệu người dùng không bao giờ rời khỏi thiết bị. Các thư viện như ONNX Runtime Web và Transformers.js dùng WebGPU để chạy các mô hình (bao gồm cả LLM nhỏ) hoàn toàn ở phía client.

Khi nào dùng Vulkan

Tình huống	Dùng Vulkan?	Tại sao / Phương án thay thế
Huấn luyện ML	Không	CUDA/Triton đơn giản và nhanh hơn trên NVIDIA
Suy luận trên GPU NVIDIA	Không	TensorRT hoặc CUDA tốt hơn
Suy luận trên GPU AMD/Intel	Có	Lựa chọn tính toán GPU duy nhất đa hãng
Suy luận mobile (Android)	Có	Vulkan là API GPU chuẩn trên Android
Suy luận mobile (iOS)	Không	Dùng Metal trực tiếp (MoltenVK thêm overhead)
Suy luận trình duyệt	WebGPU	Xây dựng trên Vulkan/Metal/DX12
Engine game + ML	Có	Engine đã dùng Vulkan cho kết xuất
Thư viện đa nền tảng	Có	Một codebase cho mọi hãng GPU
Học lập trình GPU	Có thể	CUDA dễ bắt đầu hơn; Vulkan dạy nhiều hơn

Bài tập lập trình (biên dịch với g++ -lvulkan, yêu cầu Vulkan SDK)

1. Biên dịch và chạy ví dụ cộng vector ở trên. Sửa shader để tính c[i] = a[i] * b[i] + a[i] (nhân-cộng hợp nhất) và kiểm tra kết quả.

2. Viết một compute shader áp dụng softmax cho một hàng dữ liệu dùng bộ nhớ dùng chung cho các bước reduction (max và sum). Kiểm tra với các giá trị đã biết.

// softmax.comp — biên dịch với: glslangValidator -V softmax.comp -o softmax.spv
#version 450

#define WG_SIZE 256

layout(local_size_x = WG_SIZE) in;

layout(set = 0, binding = 0) buffer Input  { float input_data[]; };
layout(set = 0, binding = 1) buffer Output { float output_data[]; };

layout(push_constant) uniform PC { uint n; };

shared float sdata[WG_SIZE];

void main() {
    uint gid = gl_GlobalInvocationID.x;
    uint lid = gl_LocalInvocationID.x;

    // Bước 1: tìm max (để ổn định số học)
    sdata[lid] = (gid < n) ? input_data[gid] : -1e30;
    barrier();
    for (uint s = WG_SIZE / 2; s > 0; s >>= 1) {
        if (lid < s) sdata[lid] = max(sdata[lid], sdata[lid + s]);
        barrier();
    }
    float maxVal = sdata[0];
    barrier();

    // Bước 2: tính exp(x - max)
    float expVal = (gid < n) ? exp(input_data[gid] - maxVal) : 0.0;
    sdata[lid] = expVal;
    barrier();

    // Bước 3: tổng các giá trị exp
    for (uint s = WG_SIZE / 2; s > 0; s >>= 1) {
        if (lid < s) sdata[lid] += sdata[lid + s];
        barrier();
    }
    float sumExp = sdata[0];

    // Bước 4: chuẩn hóa
    if (gid < n) {
        output_data[gid] = expVal / sumExp;
    }
}

1. Sửa mã host C++ để đo hiệu năng compute shader: tính thời gian dispatch (loại trừ phần thiết lập) dùng Vulkan timestamp queries hoặc fence phía CPU, và tính băng thông đạt được theo GB/s.