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 计算需要 VulkanCUDA 主导了 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 概念映射理解这套映射迁移成本几乎为零VulkanCUDA含义WorkgroupBlock可共享内存的线程组InvocationThread单个执行单元gl_GlobalInvocationIDblockIdx * blockDim threadIdx全局线程索引gl_LocalInvocationIDthreadIdxworkgroup 内线程索引gl_WorkGroupIDblockIdxworkgroup 索引local_size_xblockDim.x每 workgroup 线程数Storage bufferGlobal memory可读写的 GPU 内存Shared memory (shared)__shared__workgroup 内快速内存Push constantKernel argument小型统一数据其中#version 450表示 GLSL 4.50layout(local_size_x 256)声明每个 workgroup 256 个调用等价于 CUDA 的blockDim.x 256三个 storage buffer 通过layout(set, binding)声明为「内核参数」push constant 则承载运行期常量n。越界保护if (idx n)与 CUDA kernel 中的边界检查语义完全一致。ReLU 与共享内存sharedbarrier()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 workgroups 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 的二维 block16×16256 线程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::vectoruint32_t readSPIRV(const std::string filename) { std::ifstream file(filename, std::ios::ate | std::ios::binary); size_t fileSize file.tellg(); std::vectoruint32_t buffer(fileSize / sizeof(uint32_t)); file.seekg(0); file.read(reinterpret_castchar*(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::vectorVkPhysicalDevice 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::vectorVkQueueFamilyProperties 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。逐段拆解十步流程每步对应一个核心概念创建 Instance实例VK_API_VERSION_1_2声明应用请求的 API 版本实例是 Vulkan 应用的根对象CUDA 中无对应物驱动自动初始化。选择物理设备GPU枚举全部物理设备取第一个vkGetPhysicalDeviceProperties打印 GPU 名称——这一步相当于 CUDA 的cudaGetDeviceProperties。查找计算队列族用VK_QUEUE_COMPUTE_BIT标志遍历队列族找到支持计算队列的那个。队列族是 Vulkan 的特色同一设备可能有图形队列、计算队列、传输队列可并行提交。创建逻辑设备与队列逻辑设备VkDevice是应用的「句柄视图」队列优先级1.0f表示默认优先级。分配缓冲示例使用host-visible host-coherent内存CPU 可直接读写速度较慢但最简单生产代码通常用device-local内存GPU 专用带宽更高配合显式传输。内存类型通过memoryTypeBits掩码与propertyFlags匹配选出。填充输入缓冲vkMapMemory获得 CPU 指针写入后vkUnmapMemoryhost-coherent 保证无需手动 flush。创建计算管线SPIR-V 二进制 →VkShaderModuleVkDescriptorSetLayout描述 3 个 storage buffer 绑定stageFlags VK_SHADER_STAGE_COMPUTE_BITpush constant 范围size sizeof(uint32_t)与着色器里的uint n对齐VkPipelineLayout把描述符集与 push constant 组装成最终布局。创建描述符集从 descriptor poolmaxSets 13 个 storage buffer分配描述符集再把bufA/bufB/bufC的VkDescriptorBufferInfo通过vkUpdateDescriptorSets写入——这是着色器里binding 0/1/2与具体缓冲的连接点。录制并提交命令缓冲vkCmdBindPipeline→vkCmdBindDescriptorSets→vkCmdPushConstants→vkCmdDispatch((N 255) / 256, 1, 1)。dispatch 的 x 维 ceil(N/256)即 workgroup 数量CUDA 网格维度。fence 用于主机等待 GPU 完成vkWaitForFencesUINT64_MAX无限超时。读取结果再次 map C 缓冲验证c[i] 3.0输出ALL CORRECT。同步的两类原语着色器内的barrier()是workgroup 内同步等价__syncthreads()用于共享内存读写排序主机侧的 fence 是GPU↔CPU同步用于提交完成确认。这正是 Vulkan「只在需要时同步」哲学的体现。实践中没人手写全部样板——应封装进辅助库vk-bootstrap实例/设备创建、VMAVulkan 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(set0, binding0) buffer A { float a[]; }; layout(set0, binding1) buffer B { float b[]; }; layout(set0, binding2) 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() -recordkp::OpTensorSyncDevice({tensorA, tensorB, tensorC}) -recordkp::OpAlgoDispatch(algorithm) -recordkp::OpTensorSyncLocal({tensorC}) -eval(); // tensorC now contains [3, 3, 3, 3, 3] }可读性大幅提升kp::Manager负责实例创建、设备选择、内存分配、描述符集与命令缓冲管理OpTensorSyncDevice/OpAlgoDispatch/OpTensorSyncLocal三个算子分别对应「数据上传」「调度算法」「结果回读」record(...)-eval()对应命令录制与提交。你只需关注着色器与数据本身。WebGPU把 GPU 计算搬进浏览器WebGPU是 WebGL 的继任者为 JavaScript 提供现代 GPU 访问它建立在 VulkanLinux/Android、MetalmacOS/iOS与 DirectX 12Windows之上抽象掉平台差异。WebGPU 使用WGSLWebGPU Shading Language而非 GLSL// add.wgsl — WebGPU compute shader group(0) binding(0) varstorage, read a: arrayf32; group(0) binding(1) varstorage, read b: arrayf32; group(0) binding(2) varstorage, read_write c: arrayf32; compute workgroup_size(256) fn main(builtin(global_invocation_id) id: vec3u32) { 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场景决策表场景用 Vulkan原因 / 替代方案ML 训练否NVIDIA 上用 CUDA/Triton 更简单更快NVIDIA GPU 推理否TensorRT 或 CUDA 更好AMD/Intel GPU 推理是唯一跨厂商的 GPU compute 选项移动推理Android是Vulkan 是 Android 的标准 GPU API移动推理iOS否直接用 MetalMoltenVK 有额外开销浏览器推理WebGPU底层就是 Vulkan/Metal/DX12游戏引擎 ML是引擎渲染已用 Vulkan顺手复用跨平台库是一份代码跑通所有 GPU 厂商学习 GPU 编程看情况CUDA 上手更容易Vulkan 教得更多这套判断与仓库第 17 章 Edge Inference 的论述互为印证llama.cpp 是单文件 C 推理引擎支持 GGUF 量化Q4/Q5/Q8与 CPUAVX/NEON、Metal、CUDA、Vulkan后端是消费级硬件上运行 LLM 的首选Android 端的 Qualcomm Adreno256–1024 ALUFP16/INT8与 ARM Malitile-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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考