Vulkan Compute 与跨平台 GPU Vulkan 是唯一一个能跑在所有主流平台上的 GPU 计算 API:NVIDIA、AMD、Intel、Apple(经 MoltenVK)、Android,甚至浏览器(经 WebGPU)。本节讲 Vulkan 架构、计算流水线、用 GLSL 写计算着色器、一个完整 GPU 计算程序的 C++ 搭建、共享内存与同步、用于浏览器的 WebGPU,以及实际的 ML 推理示例。 CUDA 在 NVIDIA 硬件上统治着 ML 训练。但并非每个部署目标都有 NVIDIA GPU。一个移动应用要跑在 Qualcomm Adreno 或 ARM Mali GPU 上;一个 Web 应用要跑在浏览器里;
Vulkan 是唯一一个能跑在所有主流平台上的 GPU 计算 API:NVIDIA、AMD、Intel、Apple(经 MoltenVK)、Android,甚至浏览器(经 WebGPU)。本节讲 Vulkan 架构、计算流水线、用 GLSL 写计算着色器、一个完整 GPU 计算程序的 C++ 搭建、共享内存与同步、用于浏览器的 WebGPU,以及实际的 ML 推理示例。
CUDA 在 NVIDIA 硬件上统治着 ML 训练。但并非每个部署目标都有 NVIDIA GPU。一个移动应用要跑在 Qualcomm Adreno 或 ARM Mali GPU 上;一个 Web 应用要跑在浏览器里;一个游戏引擎需要同时支持 AMD、Intel 和 NVIDIA。对所有这些场景,Vulkan 就是答案。
Vulkan 很啰嗦——一个「hello world」计算程序大概要 300 行 C++。但这种啰嗦是显式控制的代价:你自己管理每一个 GPU 资源(内存、流水线、命令缓冲区)。这种控制带来最大性能和可移植性,代价是开发速度。
Vulkan 是由 Khronos Group(OpenGL 背后的同一个组织)创建的低层 GPU API。与 CUDA(隐藏 GPU 资源管理)不同,Vulkan 要求你显式管理:
这与 CUDA 的 cudaMalloc + 内核启动模型截然不同。CUDA 里驱动在幕后处理大部分这些事;Vulkan 里你全自己做。
Vulkan 的显式化有两个原因:
驱动简洁:OpenGL 驱动极其复杂(它得猜应用意图并据此优化)。Vulkan 把这份责任交给应用,让驱动更薄、更可预测、也更容易在各厂商之间正确实现。
性能:对内存布局、同步、命令批处理的显式控制,让应用能做出最优决策。在 CUDA 里,驱动可能插入不必要的同步;在 Vulkan 里,你只在需要时才同步。
// add.comp —— 编译:glslangValidator -V add.comp -o add.spv #version 450 // Workgroup 大小:每个 workgroup 256 个 invocation(= CUDA 里每块的线程数) layout(local_size_x = 256) in; // 缓冲区绑定(类似内核参数) 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:小的 uniform 数据(类似内核参数) layout(push_constant) uniform PushConstants { uint n; // 元素个数 }; void main() { uint idx = gl_GlobalInvocationID.x; // 全局线程索引 if (idx < n) { c[idx] = a[idx] + b[idx]; } }
| Vulkan | CUDA | 含义 |
|---|---|---|
| Workgroup | Block | 能共享内存的线程组 |
| Invocation | Thread | 单个执行单元 |
gl_GlobalInvocationID |
blockIdx * blockDim + threadIdx |
全局线程索引 |
gl_LocalInvocationID |
threadIdx |
workgroup 内的线程索引 |
gl_WorkGroupID |
blockIdx |
workgroup 索引 |
local_size_x |
blockDim.x |
每 workgroup 的线程数 |
| Storage buffer | Global memory | 可读写的 GPU 内存 |
共享内存(shared) |
__shared__ |
每 workgroup 的快速内存 |
| Push constant | 内核参数 | 小的 uniform 数据 |
// 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; }; // 共享内存(等价于 CUDA 的 __shared__) shared float tile[256]; void main() { uint gid = gl_GlobalInvocationID.x; uint lid = gl_LocalInvocationID.x; // 加载到共享内存 if (gid < n) { tile[lid] = input_data[gid]; } // 屏障:等 workgroup 内所有 invocation 加载完 barrier(); // 等价于 CUDA 的 __syncthreads() // 计算 ReLU if (gid < n) { output_data[gid] = max(tile[lid], 0.0); } }
// 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; // 加载到共享内存 sdata[lid] = (gid < n) ? input_data[gid] : 0.0; barrier(); // workgroup 内的树形归约 for (uint stride = 128; stride > 0; stride >>= 1) { if (lid < stride) { sdata[lid] += sdata[lid + stride]; } barrier(); } // 线程 0 写入这个 workgroup 的部分和 if (lid == 0) { partial_sums[wgid] = sdata[0]; } }
// 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++) { // 把 A 和 B 的一块瓦片加载到共享内存 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(); // 计算部分点积 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; } }
// vulkan_compute.cpp —— 一个最小但完整的 Vulkan 计算示例 // 编译:g++ -O3 -o vulkan_compute vulkan_compute.cpp -lvulkan // 需要:已安装 Vulkan SDK,已从 add.comp 编译出 add.spv #include <vulkan/vulkan.h> #include <iostream> #include <vector> #include <fstream> #include <cassert> // 辅助:读 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. 创建 Vulkan 实例 ========== 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. 选择物理设备(GPU) ========== uint32_t deviceCount = 0; vkEnumeratePhysicalDevices(instance, &deviceCount, nullptr); std::vector<VkPhysicalDevice> devices(deviceCount); vkEnumeratePhysicalDevices(instance, &deviceCount, devices.data()); VkPhysicalDevice physicalDevice = devices[0]; // 用第一个 GPU // 打印 GPU 名字 VkPhysicalDeviceProperties props; vkGetPhysicalDeviceProperties(physicalDevice, &props); std::cout << "Using GPU: " << props.deviceName << "\n"; // ========== 3. 找计算队列族 ========== 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. 创建逻辑设备与队列 ========== 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. 分配缓冲区(A、B、C) ========== // 为简洁起见,这里用 host-visible 内存(慢但简单) 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); // 找 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. 填充输入缓冲区 ========== 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. 创建计算流水线 ========== 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); // 描述符集布局(告诉 Vulkan 关于缓冲区绑定的信息) 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 范围 VkPushConstantRange pushRange{}; pushRange.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT; pushRange.offset = 0; pushRange.size = sizeof(uint32_t); // 流水线布局 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); // 计算流水线 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. 描述符集(把缓冲区绑定到着色器) ========== 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); // 把缓冲区引用写进描述符集 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. 录制并提交命令缓冲区 ========== 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); // 启动 workgroup vkEndCommandBuffer(cmdBuf); // 提交 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. 读取结果 ========== float* ptrC; vkMapMemory(device, memC, 0, bufferSize, 0, (void**)&ptrC); std::cout << "Results: c[0]=" << ptrC[0] << " c[1]=" << ptrC[1] << " (expected 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 ? "ALL CORRECT" : "ERRORS FOUND") << "\n"; vkUnmapMemory(device, memC); // ========== 清理(略) ========== 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; }
是的,做一次向量加法要约 200 行。 对比 CUDA 大约 30 行。这就是显式化的代价。但请注意:每一行都有它的目的。没有隐藏的驱动决策、没有隐式同步、没有意外的分配。一切由你掌控。
实际中你会把这些样板包进一个辅助库(或者用现成的,比如 vk-bootstrap、做内存分配的 VMA,或面向 ML 的 Vulkan 计算 kompute)。
#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 现在是 [3, 3, 3, 3, 3] }
WebGPU 是 WebGL 的继任者,从 JavaScript 提供现代 GPU 访问。它建立在 Vulkan(Linux/Android)、Metal(macOS/iOS)和 DirectX 12(Windows)之上,抹平了平台差异。
WebGPU 用 WGSL(WebGPU Shading Language)而不是 GLSL:
// add.wgsl —— WebGPU 计算着色器 @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]; }
const adapter = await navigator.gpu.requestAdapter(); const device = await adapter.requestDevice(); // 创建缓冲区 const bufferA = device.createBuffer({ size: N * 4, usage: GPUBufferUsage.STORAGE, mappedAtCreation: true }); new Float32Array(bufferA.getMappedRange()).fill(1.0); bufferA.unmap(); // ...(B 和 C 类似) // 用 WGSL 着色器创建流水线 const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code: wgslSource }), entryPoint: 'main' } }); // 派发 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()]);
| 场景 | 用 Vulkan? | 原因 / 替代方案 |
|---|---|---|
| ML 训练 | 否 | CUDA/Triton 在 NVIDIA 上更简单更快 |
| NVIDIA GPU 上推理 | 否 | TensorRT 或 CUDA 更好 |
| AMD/Intel GPU 上推理 | 是 | 唯一的跨厂商 GPU 计算选项 |
| 移动端推理(Android) | 是 | Vulkan 是 Android 上的标准 GPU API |
| 移动端推理(iOS) | 否 | 直接用 Metal(MoltenVK 有额外开销) |
| 浏览器推理 | WebGPU | 建立在 Vulkan/Metal/DX12 之上 |
| 游戏引擎 + ML | 是 | 引擎本来就用 Vulkan 做渲染 |
| 跨平台库 | 是 | 一套代码覆盖所有 GPU 厂商 |
| 学习 GPU 编程 | 也许 | CUDA 更易入门;Vulkan 教得更深 |
编译并运行上面的向量加法示例。把着色器改成计算 c[i] = a[i] * b[i] + a[i](融合乘加)并验证结果。
写一个计算着色器,用共享内存做归约步骤(最大值和求和),对一行数据做 softmax。用已知值做测试。
// softmax.comp —— 编译: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; // 第 1 步:找最大值(保证数值稳定) 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(); // 第 2 步:计算 exp(x - max) float expVal = (gid < n) ? exp(input_data[gid] - maxVal) : 0.0; sdata[lid] = expVal; barrier(); // 第 3 步:求 exp 值之和 for (uint s = WG_SIZE / 2; s > 0; s >>= 1) { if (lid < s) sdata[lid] += sdata[lid + s]; barrier(); } float sumExp = sdata[0]; // 第 4 步:归一化 if (gid < n) { output_data[gid] = expVal / sumExp; } }