企业官网建设流程全解析

Maths-CS-AI Compendium 精讲：Vulkan Compute 与跨平台 GPU 计算——从 GLSL 着色器到 C++/WebGPU 的 ML 推理实战

【免费下载链接】maths-cs-ai-compendiumBecome a cracked AI/ML researcher/engineer with this unconventional textbook covering maths, computing, and ML with intuition.项目地址: https://gitcode.com/GitHub_Trending/mat/maths-cs-ai-compendium

本指南是《Maths, CS & AI Compendium》第 16 章「SIMD 与 GPU 编程」的收官篇（Vulkan Compute and Cross-Platform GPU），围绕 Vulkan 的架构与计算管线展开：如何用 GLSL 编写计算着色器、如何用 C++ 搭建一个完整的 GPU 计算程序、共享内存与同步原语如何工作，以及 WebGPU 如何在浏览器中复用同一套思想。读完本文，你将掌握一套与 CUDA 概念一一映射的跨厂商 GPU 编程方法，并能在 AMD/Intel/ARM 移动 GPU 甚至浏览器中落地 ML 推理（如 llama.cpp 的 Vulkan 后端）。

为什么跨平台 GPU 计算需要 Vulkan

CUDA 主导了 NVIDIA 硬件上的 ML 训练，但并非每个部署目标都有一块 NVIDIA GPU：手机 App 跑在 Qualcomm Adreno 或 ARM Mali 上，Web 应用只能依赖浏览器运行时，游戏引擎必须同时兼容 AMD、Intel 与 NVIDIA。对这些场景，Vulkan是通用答案——它是唯一能在 NVIDIA、AMD、Intel、Apple（经 MoltenVK）、Android 乃至浏览器（经 WebGPU）上运行的主流 GPU compute API。

代价是冗长：一个「hello world」计算程序大约需要 300 行 C++。但这份冗长换来的正是显式控制——实例与设备、内存、缓冲、描述符集、计算管线、命令缓冲、队列提交、同步，全部由你亲自管理。这与 CUDA 的cudaMalloc+ kernel launch 模型截然不同：CUDA 驱动替你处理了绝大多数资源管理，而 Vulkan 把决定权交还给你，从而获得最大性能与可移植性。

为什么 Vulkan 如此冗长

• 驱动更简单：OpenGL 驱动极其复杂，必须猜测应用意图并据此优化。Vulkan 把这份责任移交给应用，驱动因此更薄、更可预测、更容易在各厂商间正确实现。
• 性能更优：显式控制内存布局、同步与命令批处理，让应用做出最优决策。CUDA 驱动可能插入不必要的同步；Vulkan 只在需要时同步。

这与仓库中 GPU Architecture and CUDA 一文的视角互为印证：CUDA 适合在单一厂商栈上快速起步，而 Vulkan 面向「一份代码跑遍所有 GPU」的跨平台诉求。第 16 章 Triton、TPUs 与 Pallas 的总结表也把 Vulkan 列为「跨平台推理」的首选：Cross-platform inference → Vulkan or ONNX Runtime → Runs on any GPU vendor。

GLSL 计算着色器：从零到分块矩阵乘

计算着色器（compute shader）是跑在 GPU 上的程序，与 CUDA kernel 等价，用GLSL编写并经glslangValidator编译为SPIR-V字节码（可移植二进制格式）。下面是仓库文档中的完整示例。

向量加法：第一个 compute shader

// add.comp — compile with: glslangValidator -V add.comp -o add.spv #version 450 // Workgroup size: 256 invocations per workgroup (= threads per block in CUDA) layout(local_size_x = 256) in; // Buffer bindings (like kernel arguments) 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: small uniform data (like a kernel parameter) layout(push_constant) uniform PushConstants { uint n; // number of elements }; void main() { uint idx = gl_GlobalInvocationID.x; // global thread index if (idx < n) { c[idx] = a[idx] + b[idx]; } }

Vulkan ↔ CUDA 概念映射（理解这套映射，迁移成本几乎为零）：

其中#version 450表示 GLSL 4.50，layout(local_size_x = 256)声明每个 workgroup 256 个调用（等价于 CUDA 的blockDim.x = 256）；三个 storage buffer 通过layout(set, binding)声明为「内核参数」，push constant 则承载运行期常量n。越界保护if (idx < n)与 CUDA kernel 中的边界检查语义完全一致。

ReLU 与共享内存：shared+barrier()

ReLU 本身是逐元素操作，并不真正需要共享内存——但这个示例展示了 GPU 编程的核心模式：load → barrier → compute → store。

// 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; }; // Shared memory (equivalent to CUDA __shared__) shared float tile[256]; void main() { uint gid = gl_GlobalInvocationID.x; uint lid = gl_LocalInvocationID.x; // Load into shared memory if (gid < n) { tile[lid] = input_data[gid]; } // Barrier: wait for all invocations in workgroup to finish loading barrier(); // equivalent to CUDA __syncthreads() // Compute ReLU if (gid < n) { output_data[gid] = max(tile[lid], 0.0); } }

shared float tile[256]对应 CUDA 的__shared__ float tile[256]，barrier()对应__syncthreads()。对于需要读取相邻线程数据的操作（卷积、归约、softmax），共享内存是性能关键。仓库第 16 章 GPU Architecture and CUDA 指出：共享内存是「由程序员管理的、block 内所有线程共享的高速缓存」，是写出快速 CUDA kernel 的关键——tiling 模式贯穿 GPU 编程，这一论断同样适用于 Vulkan。

并行归约（求和）：经典树形归约

// 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; // Load into shared memory sdata[lid] = (gid < n) ? input_data[gid] : 0.0; barrier(); // Tree reduction within the workgroup for (uint stride = 128; stride > 0; stride >>= 1) { if (lid < stride) { sdata[lid] += sdata[lid + stride]; } barrier(); } // Thread 0 writes the workgroup's partial sum if (lid == 0) { partial_sums[wgid] = sdata[0]; } }

这是与 CUDA 完全相同的经典归约模式：每个 workgroup 产出一个部分和，第二次 dispatch 再把这些部分和归约为最终结果。树形归约每步活跃线程减半：256 → 128 → 64 → … → 1，注意越界线程以0.0填充，保证归约正确。

分块矩阵乘法：GLSL 版 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++) { // Load tile of A and B into shared memory 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(); // Compute partial dot product 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; } }

这里local_size_x/y = 16对应 CUDA 的二维 block（16×16=256 线程），tileA/tileB对应 CUDA 的__shared__ float tile_A[TILE_SIZE][TILE_SIZE]。算法与 GPU Architecture and CUDA 第 182 行起的matmul_tiledkernel 逐行同构：按 t 轮遍历 K 维度上的所有 tile，先加载分块到共享内存、barrier()，再计算局部点积、barrier()防止下一轮加载覆盖未消费数据。tiling 为什么有效：无分块时，每个线程每次乘法都要访问全局内存；分块后一个 TILE_SIZE×TILE_SIZE 的数据块只从全局内存加载一次，被 block 内所有线程复用，复用因子为 TILE_SIZE，从而把全局内存流量降低相应倍数。边界处补零（? A[...] : 0.0）保证 K 不能被 TILE_SIZE 整除时结果仍正确。

完整 C++ Vulkan 计算程序：10 步走完 GPU 管线

着色器只是「简单」的部分；真正繁琐的是 C++ 样板代码：创建实例、分配内存、绑定缓冲、提交命令。仓库文档给出了一份最小但完整的示例，其十步流程正是 Vulkan 计算管线的标准骨架：

// vulkan_compute.cpp — a minimal but complete Vulkan compute example // Compile: g++ -O3 -o vulkan_compute vulkan_compute.cpp -lvulkan // Requires: Vulkan SDK installed, add.spv compiled from add.comp #include <vulkan/vulkan.h> #include <iostream> #include <vector> #include <fstream> #include <cassert> // Helper: read SPIR-V file 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. Create 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. Select 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]; // use first GPU // Print GPU name VkPhysicalDeviceProperties props; vkGetPhysicalDeviceProperties(physicalDevice, &props); std::cout << "Using GPU: " << props.deviceName << "\n"; // ========== 3. Find 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. Create Logical Device and 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. Allocate Buffers (A, B, C) ========== // For brevity, this uses host-visible memory (slower but simpler) auto createBuffer = & { 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); // Find host-visible memory type 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. Fill Input Buffers ========== 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. Create 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 (tells Vulkan about the buffer bindings) 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 (bind buffers to 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); // Write buffer references into the 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. Record and Submit 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); // launch workgroups vkEndCommandBuffer(cmdBuf); // Submit 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. Read Results ========== 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); // ========== Cleanup (abbreviated) ========== 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 行 C++，而 CUDA 只需约 30 行。但请注意：每一行都有明确目的，没有隐藏的驱动决策、隐式同步或意外分配——一切尽在掌握。编译命令为g++ -O3 -o vulkan_compute vulkan_compute.cpp -lvulkan，前提是已安装 Vulkan SDK 且已用glslangValidator -V add.comp -o add.spv生成add.spv。

逐段拆解十步流程，每步对应一个核心概念：

1. 创建 Instance（实例）：VK_API_VERSION_1_2声明应用请求的 API 版本；实例是 Vulkan 应用的根对象，CUDA 中无对应物（驱动自动初始化）。
2. 选择物理设备（GPU）：枚举全部物理设备，取第一个；vkGetPhysicalDeviceProperties打印 GPU 名称——这一步相当于 CUDA 的cudaGetDeviceProperties。
3. 查找计算队列族：用VK_QUEUE_COMPUTE_BIT标志遍历队列族，找到支持计算队列的那个。队列族是 Vulkan 的特色：同一设备可能有图形队列、计算队列、传输队列，可并行提交。
4. 创建逻辑设备与队列：逻辑设备（VkDevice）是应用的「句柄视图」，队列优先级1.0f表示默认优先级。
5. 分配缓冲：示例使用host-visible + host-coherent内存（CPU 可直接读写，速度较慢但最简单）；生产代码通常用device-local内存（GPU 专用，带宽更高）配合显式传输。内存类型通过memoryTypeBits掩码与propertyFlags匹配选出。
6. 填充输入缓冲：vkMapMemory获得 CPU 指针，写入后vkUnmapMemory；host-coherent 保证无需手动 flush。
7. 创建计算管线：SPIR-V 二进制 →VkShaderModule；VkDescriptorSetLayout描述 3 个 storage buffer 绑定（stageFlags = VK_SHADER_STAGE_COMPUTE_BIT）；push constant 范围size = sizeof(uint32_t)与着色器里的uint n对齐；VkPipelineLayout把描述符集与 push constant 组装成最终布局。
8. 创建描述符集：从 descriptor pool（maxSets = 1，3 个 storage buffer）分配描述符集，再把bufA/bufB/bufC的VkDescriptorBufferInfo通过vkUpdateDescriptorSets写入——这是着色器里binding = 0/1/2与具体缓冲的连接点。
9. 录制并提交命令缓冲：vkCmdBindPipeline→vkCmdBindDescriptorSets→vkCmdPushConstants→vkCmdDispatch((N + 255) / 256, 1, 1)。dispatch 的 x 维 =ceil(N/256)，即 workgroup 数量（CUDA 网格维度）。fence 用于主机等待 GPU 完成（vkWaitForFences+UINT64_MAX无限超时）。
10. 读取结果：再次 map C 缓冲，验证c[i] = 3.0，输出ALL CORRECT。

同步的两类原语：着色器内的barrier()是workgroup 内同步（等价__syncthreads()），用于共享内存读写排序；主机侧的 fence 是GPU↔CPU同步，用于提交完成确认。这正是 Vulkan「只在需要时同步」哲学的体现。

实践中，没人手写全部样板——应封装进辅助库：vk-bootstrap（实例/设备创建）、VMA（Vulkan Memory Allocator，内存分配）、Kompute（面向 ML 的 Vulkan compute 封装）。

Kompute：面向 ML 的 Vulkan 简化封装

Kompute是开源 C++ 库，把 Vulkan 样板代码打包为易用接口。同样的向量加法变成：

#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 now contains [3, 3, 3, 3, 3] }

可读性大幅提升：kp::Manager负责实例创建、设备选择、内存分配、描述符集与命令缓冲管理，OpTensorSyncDevice/OpAlgoDispatch/OpTensorSyncLocal三个算子分别对应「数据上传」「调度算法」「结果回读」，record(...)->eval()对应命令录制与提交。你只需关注着色器与数据本身。

WebGPU：把 GPU 计算搬进浏览器

WebGPU是 WebGL 的继任者，为 JavaScript 提供现代 GPU 访问；它建立在 Vulkan（Linux/Android）、Metal（macOS/iOS）与 DirectX 12（Windows）之上，抽象掉平台差异。WebGPU 使用WGSL（WebGPU Shading Language）而非 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]; }

WGSL 与 GLSL 的映射清晰可见：@group(0) @binding(i)↔layout(set = 0, binding = i)，@workgroup_size(256)↔layout(local_size_x = 256)，@builtin(global_invocation_id)↔gl_GlobalInvocationID。JavaScript 侧（浓缩版）：

const adapter = await navigator.gpu.requestAdapter(); const device = await adapter.requestDevice(); // Create buffers const bufferA = device.createBuffer({ size: N * 4, usage: GPUBufferUsage.STORAGE, mappedAtCreation: true }); new Float32Array(bufferA.getMappedRange()).fill(1.0); bufferA.unmap(); // ... (similar for B and C) // Create pipeline from WGSL shader 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()]);

WebGPU 对 ML 的意义：浏览器内推理意味着零服务器成本、零网络延迟，且用户数据不出设备。ONNX Runtime Web、Transformers.js 等库正是用 WebGPU 在客户端完整运行模型（包括小型 LLM）。

何时使用 Vulkan：场景决策表

这套判断与仓库第 17 章 Edge Inference 的论述互为印证：llama.cpp 是单文件 C++ 推理引擎，支持 GGUF 量化（Q4/Q5/Q8）与 CPU（AVX/NEON）、Metal、CUDA、Vulkan后端，是消费级硬件上运行 LLM 的首选；Android 端的 Qualcomm Adreno（256–1024 ALU，FP16/INT8）与 ARM Mali（tile-based 架构，影响内存访问模式）均通过 Vulkan compute 暴露算力。第 17 章 Scaling and Deployment 的推理引擎对比表同样将 llama.cpp 标记为CPU/Metal/CUDA/Vulkan, GGUF quantisation, portable——可见 Vulkan 是「一份引擎跑遍所有硬件」的关键拼图。

动手练习：从加法到 softmax 与带宽基准

以下练习均以g++ -lvulkan编译，需先安装 Vulkan SDK。

练习 1：融合乘加（FMA）

编译并运行上面的向量加法示例，把着色器改为c[i] = a[i] * b[i] + a[i]（融合乘加），验证结果。

练习 2：共享内存版 softmax

为一行数据编写 softmax 着色器，用共享内存完成 max 与 sum 两趟归约，保证数值稳定性：

// softmax.comp — compile with: 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; // Step 1: find max (for numerical stability) 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(); // Step 2: compute exp(x - max) float expVal = (gid < n) ? exp(input_data[gid] - maxVal) : 0.0; sdata[lid] = expVal; barrier(); // Step 3: sum of exp values for (uint s = WG_SIZE / 2; s > 0; s >>= 1) { if (lid < s) sdata[lid] += sdata[lid + s]; barrier(); } float sumExp = sdata[0]; // Step 4: normalise if (gid < n) { output_data[gid] = expVal / sumExp; } }

这段代码把本文所有核心技巧串成一条线：载入共享内存 → 归约求 max →exp(x - max)保证数值稳定 → 归约求和 → 归一化。其中-1e30作为 -∞ 哨兵值参与 max 归约，越界线程在求和阶段贡献0.0。

练习 3：带宽基准测试

修改 C++ 宿主代码，用 Vulkan timestamp query（时间戳查询）或 CPU 侧 fence 计时 dispatch（排除初始化），计算实际带宽 GB/s =3 * N * 4 字节 / 耗时。对照第 16 章 GPU Architecture and CUDA 中 CUDA 版matmul_tiled的cudaEventElapsedTime基准做法，可以直观比较同一算法在 Vulkan 与 CUDA 下的差距。

小结

Vulkan 以显式控制换来了真正的跨平台：从 GLSL 着色器里的shared/barrier()到 C++ 侧的描述符集与命令缓冲，再到浏览器中的 WebGPU，核心心智模型与 CUDA 完全同构——workgroup 即 block、invocation 即 thread、tiling 即性能。对 ML 工程而言，当部署目标是 AMD/Intel 桌面 GPU、Android 移动端或浏览器时，Vulkan/WebGPU 几乎是唯一「一份代码通吃」的路线；这也正是本仓库把 Triton/TPU（第 16 章第 5 节）、边缘推理 与本文件编排进同一学习路径的原因：先懂硬件（01 硬件基础、04 CUDA），再懂抽象（Triton/Vulkan/WebGPU），最后落到部署决策（第 17 章）。若你想快速上手实践，直接在本仓库根目录下找到本文件，把四个 GLSL 着色器与 C++ 宿主代码原样编译运行，就是最完整的「hello world」进阶路线。

【免费下载链接】maths-cs-ai-compendiumBecome a cracked AI/ML researcher/engineer with this unconventional textbook covering maths, computing, and ML with intuition.项目地址: https://gitcode.com/GitHub_Trending/mat/maths-cs-ai-compendium