052、从GPU Dialect到LLVM/SPIR-V的降级上周调试一个自定义算子时，发现kernel在NVIDIA卡上跑得好好的，换到Intel GPU上直接崩了。报错信息指向SPIR-V的workgroup size对齐问题。翻了一整天MLIR的降级流水线，才意识到GPU Dialect到LLVM/SPIR-V的降级路径里藏着不少“暗坑”。今天把这块的笔记整理出来，权当给自己留个备忘。降级路径的骨架MLIR里GPU Dialect的定位很明确——它是一层硬件无关的并行计算抽象。你写gpu.launch、gpu.thread_id，它不关心底层是CUDA还是OpenCL还是Vulkan。真正干活的是降级pass：把GPU Dialect拆解成LLVM Dialect（走NVVM路径）或者SPIR-V Dialect（走通用GPU路径）。我习惯把降级分成三个阶段：第一阶段：GPU Dialect → GPUToLLVMConversionPass这个pass把gpu.launch展开成循环和函数调用。比如gpu.thread_id会被替换成llvm.call @__gpu_thread_id_x()。注意这里有个坑——不同GPU后端对线程ID的编号方式不同。NVIDIA用threadIdx.x，AMD用local_id.x，SPIR-V用