企业官网建设流程全解析

Slang 地址空间（Address Spaces）全面解析：11 种内存视图、跨后端映射与一致性测试验证

【免费下载链接】slangMaking it easier to work with shaders项目地址: https://gitcode.com/GitHub_Trending/sl/slang

导读

本文以 Slang 语言参考手册中的 basics-memory-model-address-spaces.md 为骨架，结合仓库中为其专门生成的一致性测试束 conformance/basics-memory-model-address-spaces，系统讲解 Slang 的 11 种地址空间：每种地址空间对应的 Slang 构造、实例作用域（Instance scope），以及它们分别如何被发射到 SPIR-V、Metal、WGSL、GLSL、HLSL 与 CUDA 六种后端。读完本文，你将掌握地址空间的选择规则、跨后端存储类映射表，以及如何通过spirv-asm/metal/wgsl等发射测试验证这些映射。

一、什么是地址空间：变量内存归属的决定规则

地址空间（Address Space）定义了一种查看内存的方式，它决定内存对象如何被线程寻址和访问。在 Slang 中，一个变量的地址空间由三方面共同决定（对应一致性测试束中的C0声明）：

• 类型（type）：例如ConstantBuffer<T>天然落入 Uniform 地址空间；
• 修饰符（modifiers）：例如static groupshared声明的是组共享地址空间；
• 属性（attributes）：例如[vk::push_constant]、[vk::specialization_constant]分别把变量送入 Push constant 与 Specialization constant 地址空间。

同一个着色器里可以同时存在多种地址空间的变量，例如一个计算着色器可以同时拥有ConstantBuffer<Params>（Uniform）、RWStructuredBuffer<float>（Storage buffer）和static groupshared float[32]（Group-shared）——这一"多地址空间共存"正是测试 address-space-by-type-spirv.slang 的验证目标。

二、11 种地址空间总览表

语言参考文档以一张规范性表格定义全部 11 种地址空间，下面完整继承该表：

文档还附带三条备注（Remarks），它们同样被一致性测试束收录为可验证声明：

• Remark 1：与图形管线阶段相关的地址空间未在此枚举，参见 Graphics Shaders and Compute Kernels。
• Remark 2：一个地址空间中的内存指针通常不能与另一个地址空间中的指针互换；特别是组共享内存的指针在线程组之间不可互换。
• Remark 3：Slang 中的地址空间与SPIR-V 存储类（storage class）大致等价——这一条是整个一致性测试束"发射优先"（emission-first）验证策略的基石。

三、逐地址空间深解：Slang 构造与六后端发射映射

一致性测试束为每个地址空间编写了至少一个.slang测试文件，每个文件同时用//TEST:SIMPLE指令把同一份源码分别发射到spirv-asm、metal、wgsl、glsl、hlsl（组共享还额外覆盖cuda），用filecheck断言各后端的存储类/存储限定符，从而把语言参考文档中的每条声明"钉死"在六个后端上。测试束的每目标回归覆盖面因此从 4 个后端扩展到了 6 个（新增 Metal 与 WGSL）。

3.1 Uniform：常量缓冲与uniform关键字

Uniform 地址空间对应两个 Slang 构造：

• C1：ConstantBuffer<T>。测试 uniform-constant-buffer-spirv.slang 验证其跨后端映射：

//TEST:SIMPLE(filecheck=CHECK):-target spirv-asm -entry computeMain -stage compute //TEST:SIMPLE(filecheck=METAL):-target metal -entry computeMain -stage compute //TEST:SIMPLE(filecheck=WGSL):-target wgsl -entry computeMain -stage compute //TEST:SIMPLE(filecheck=GLSL):-target glsl -entry computeMain -stage compute //TEST:SIMPLE(filecheck=HLSL):-target hlsl -entry computeMain -stage compute struct Params { float scale; int offset; }; ConstantBuffer<Params> cb; [numthreads(1, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<float> output) { output[0] = cb.scale + float(cb.offset); } // CHECK: OpVariable {{.*}} Uniform // METAL: constant* cb_ // WGSL: var<uniform> cb_ // GLSL: layout(std140) uniform block_Params_ // HLSL: cbuffer cb_{{.*}} : register(b

映射结论：SPIR-V 发射Uniform存储类；Metal 为constant*指针；WGSL 为var<uniform>；GLSL 为layout(std140)uniform block；HLSL 为b寄存器上的cbuffer。

• C2：入口点参数上的uniform关键字。测试 uniform-keyword-spirv.slang 表明它同样属于 Uniform 地址空间，但值得注意的细节是：在 SPIR-V 中uniform入口点参数被发射进PushConstant块，Metal 为constant*，WGSL 为var<uniform>——这与ConstantBuffer<T>恒映射到Uniform存储类并不相同（该差异也进入了"文档缺口"清单，见第六节）。

3.2 Image：纹理类型

C3：Texture2D<T>（以及Texture1D<T>等其它纹理类型）属于 Image 地址空间，实例作用域为所有线程。测试 image-texture2d-spirv.slang 与 image-texture1d-spirv.slang 验证其映射：SPIR-V 发射UniformConstant存储类，Metal 为texture2d<>，WGSL 为texture_2d<>。

3.3 Push constant：高频更新的小型常量

C4：[vk::push_constant]标注的结构体属于 Push constant 地址空间。测试 push-constant-spirv.slang 验证映射：

struct PushData { float scale; int count; }; [[vk::push_constant]] PushData pushConst; [numthreads(1, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<float> output) { output[0] = pushConst.scale * float(pushConst.count); } // CHECK: %pushConst = OpVariable {{.*}} PushConstant // METAL: constant* pushConst_ // WGSL: var<uniform> pushConst_ // GLSL: layout(push_constant) // HLSL: cbuffer pushConst_{{.*}} : register(b

映射结论：SPIR-V 发射PushConstant存储类；由于 Metal/WGSL/HLSL 没有原生 push constant 概念，Slang 分别将其降级为constant*缓冲参数、var<uniform>与b寄存器上的cbuffer；GLSL 则保留layout(push_constant)block。

C5（边界用例）：不带vk前缀的[push_constant]同样产生PushConstant存储类，验证于 push-constant-attr-spirv.slang，Metal/WGSL 降级行为与带前缀形式一致。

3.4 Storage buffer：只读与可读写结构化缓冲

C6/C7：RWStructuredBuffer<T>（可读写）与StructuredBuffer<T>（只读）都属于 Storage buffer 地址空间。测试 storage-buffer-rw-spirv.slang 完整展示了只读/可读写两种形态的区分：

RWStructuredBuffer<float> rwbuf; StructuredBuffer<float> robuf; [numthreads(1, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<float> output) { float v = robuf[0]; rwbuf[0] = v + 1.0; output[0] = rwbuf[0]; } // CHECK-DAG: %rwbuf = OpVariable {{.*}} StorageBuffer // CHECK-DAG: %robuf = OpVariable {{.*}} StorageBuffer {{.*}} NonWritable // METAL: float device* robuf_ // WGSL: var<storage, read> robuf_ // WGSL-DAG: var<storage, read_write> rwbuf_ // GLSL: readonly buffer StructuredBuffer_float // HLSL: StructuredBuffer<float{{.*}}> robuf_{{.*}} : register(t // HLSL: RWStructuredBuffer<float{{.*}}> rwbuf_{{.*}} : register(u

映射结论：两者在 SPIR-V 中都是StorageBuffer，但只读形态额外带NonWritable装饰；Metal 均为device*；WGSL 区分var<storage, read>与var<storage, read_write>；GLSL 区分readonly buffer与普通buffer（std430）；HLSL 按访问方式拆分为 SRV（t寄存器）与 UAV（u寄存器）。HLSL 侧的等价验证见 storage-buffer-hlsl.slang。

3.5 Group-shared：线程组内共享内存

C8：全局作用域的static groupshared属于 Group-shared 地址空间，实例作用域为线程组。这是覆盖后端最全的一组测试（4 个文件、6 个后端）：

• groupshared-workgroup-spirv.slang：SPIR-V 发射Workgroup存储类；

static groupshared float sharedData[32]; [numthreads(32, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<float> output) { sharedData[tid.x] = float(tid.x); GroupMemoryBarrierWithGroupSync(); output[tid.x] = sharedData[31u - tid.x]; } // CHECK: %sharedData = OpVariable {{.*}} Workgroup // METAL: threadgroup {{.*}} sharedData{{.*}} // WGSL: var<workgroup> sharedData{{.*}} // GLSL: shared float{{.*}} sharedData_ // HLSL: static groupshared float{{.*}} sharedData_

• groupshared-workgroup-hlsl.slang：HLSL 保留groupshared限定符；
• groupshared-workgroup-glsl.slang：GLSL 发射shared存储限定符；
• groupshared-cuda.slang：CUDA 后端发射__shared__；此外 Metal 为threadgroup，WGSL 为var<workgroup>。

C9（运行时语义声明）：组共享变量实例由线程组共享——同一组内的线程在GroupMemoryBarrierWithGroupSync()屏障之后可以互相读取对方写入的值。这条声明被列入"未测试声明"，原因见第五节。

3.6 Function：函数参数与局部变量

C10：函数参数与非 static 局部变量属于 Function 地址空间，实例作用域为线程，仅对本次函数调用可见。测试 function-local-var-spirv.slang 刻意使用带inout参数的Counter结构体局部变量，使其无法被 SSA 折叠，从而在 SPIR-V 中保留可断言的OpVariable：

struct Counter { int val; }; [__NoInline] int countUp(inout Counter c, int n) { for (int i = 0; i < n; i++) c.val++; return c.val; } [numthreads(1, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<int> output) { Counter c; c.val = 0; output[0] = countUp(c, int(tid.x) + 1); } // CHECK: OpVariable {{.*}} Function // METAL: thread Counter_ // WGSL: ptr<function, Counter_

映射结论：SPIR-V 发射Function存储类；Metal 为thread限定符的结构体局部变量；WGSL 以ptr<function, ...>传参；GLSL/HLSL 在函数体内发射无存储限定符的局部声明（无限定符即 Function 空间）。

3.7 Thread-local：每个线程一个实例的全局static

C11：全局作用域、非groupshared的static变量属于 Thread-local 地址空间，每个线程获得独立实例。测试 thread-local-static-hlsl.slang 用[noinline]辅助函数跨函数读写该变量，保证它不会被常量折叠：

static int g_perThread; [noinline] int readAndSet(int v) { int old = g_perThread; g_perThread = v; return old; } // CHECK: static int g_perThread{{.*}} // METAL: thread KernelContext // WGSL: var<private> g_perThread{{.*}} // GLSL: int g_perThread_ // SPIRV: %g_perThread = OpVariable {{.*}} Private

映射结论：HLSL 保留模块作用域的static；Metal 内联进带thread地址空间的KernelContext结构；WGSL 发射var<private>（private 地址空间即每线程）；GLSL 发射无存储限定符的全局变量（GLSL 全局变量默认逐调用实例，shared才是组作用域）；SPIR-V 发射Private存储类。

3.8 Input 与 Output：图形阶段的入口数据

C12：非 uniform 的入口点输入参数属于 Input 地址空间，实例作用域为线程。测试 input-address-space-spirv.slang 以SV_DispatchThreadID为例：

[numthreads(4, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<uint> output) { output[tid.x] = tid.x; } // CHECK: OpVariable {{.*}} Input // CHECK-DAG: BuiltIn GlobalInvocationId // METAL: thread_position_in_grid // WGSL: @builtin(global_invocation_id) // GLSL: gl_GlobalInvocationID // HLSL: void computeMain(uint3 tid_{{.*}} : SV_DispatchThreadID)

映射结论：SPIR-V 发射Input存储类并带BuiltIn GlobalInvocationId装饰；Metal 为thread_position_in_grid内建属性；WGSL 为@builtin(global_invocation_id)；GLSL 为gl_GlobalInvocationID；HLSL 保持SV_DispatchThreadID语义参数。

C13：入口点输出参数与返回值属于 Output 地址空间。测试 output-address-space-spirv.slang 用一个返回SV_Position与TEXCOORD0的顶点着色器验证：

struct VSOutput { float4 pos : SV_Position; float2 uv : TEXCOORD0; }; VSOutput vsMain(float3 pos : POSITION, float2 uv : TEXCOORD0) { VSOutput o; o.pos = float4(pos, 1.0); o.uv = uv; return o; } // CHECK: OpVariable {{.*}} Output // METAL-DAG: vertex // METAL-DAG: position // WGSL-DAG: @vertex // WGSL-DAG: @builtin(position) // GLSL: out vec2 entryPointParam_vsMain_uv_ // GLSL: gl_Position = // HLSL: float4 pos_{{.*}} : SV_Position;

映射结论：SPIR-V 发射Output存储类；Metal 以[[vertex]]入口点配合[[position]]输出字段；WGSL 以@vertex配合@builtin(position)；GLSL 以outvarying 与gl_Position内建变量；HLSL 保持语义标注的返回结构体字段。

3.9 Specialization constant：管线创建时固定的常量

C14/C15：[vk::specialization_constant]（以及不带前缀的[SpecializationConstant]）标注的变量属于 Specialization constant 地址空间。与 Uniform 变量不同，特化常量发射为OpSpecConstant（而非OpVariable），每个都带SpecId装饰。测试 spec-constant-spirv.slang 用CHECK-DAG同时断言两种属性拼写：

[[vk::specialization_constant]] const int WORK_SIZE = 4; [[SpecializationConstant]] const float SCALE_FACTOR = 1.5; [numthreads(1, 1, 1)] void computeMain(uint3 tid : SV_DispatchThreadID, RWStructuredBuffer<float> output) { output[0] = float(WORK_SIZE) * SCALE_FACTOR; } // CHECK-DAG: OpSpecConstant %int {{.*}} // CHECK-DAG: OpSpecConstant %float {{.*}} // CHECK-DAG: OpDecorate {{.*}} SpecId // METAL: function_constant( // METAL-DAG: is_function_constant_defined( // WGSL: @id( // WGSL-DAG: override WORK_SIZE_ // GLSL: layout(constant_id = {{[0-9]+}}) // GLSL-NEXT: const int WORK_SIZE_

映射结论：SPIR-V 为OpSpecConstant+SpecId装饰；Metal 为function_constant（带回退判断）；WGSL 为带@id的override；GLSL 为layout(constant_id = N) const。非 vk 形式的等价验证见 spec-constant-attr-spirv.slang。

3.10 Host：Slang 程序通常无法直接访问

C16：Host 地址空间没有对应的 Slang 构造，仅由宿主程序使用。它通常是应用进程的虚拟地址空间，Slang 程序通常无法直接访问——唯一的例外是编译为 C++ 目标时。由于该声明不涉及static groupshared等特殊构造，其验证被划归到 types-pointer 测试束（host-cpp目标行为），不在本束内覆盖。

四、一致性测试束：从文档声明到可执行断言

该束由claude-sonnet-4-6自动生成（generated: true），每个测试文件头部带有//META元数据，记录其对应的文档锚点（doc_ref=docs/language-reference/basics-memory-model-address-spaces.md#address-spaces）、purpose、intent（functional / boundary）与pipeline_stage（emit）。核心的测试机制是：

1. 每条//TEST:SIMPLE(filecheck=XXX):-target <target> -entry <entry> -stage <stage>指令把同一份.slang源码发射到指定后端；
2. filecheck=CHECK/METAL/WGSL/GLSL/HLSL/SPIRV分别匹配各后端的输出片段；
3. 所有测试的intent均为functional，pipeline_stage均为emit——即验证"发射输出正确"而非"运行结果正确"。

README 中"功能覆盖"（Functional coverage）表把 20 条声明（C0–C19）与具体测试文件一一对应，下表完整继承：

五、未测试声明：哪些声明无法用发射测试验证

README 的## Untested claims表诚实记录了 6 条无法在本束内验证的声明及其原因，完整继承如下：

六、文档缺口观察：一致性测试反哺语言参考

测试束在验证过程中发现的三处文档问题，被记录在## Doc gaps observed表中，同样是本主题下值得开发者注意的坑点：

1. uniform关键字映射与文档表述不一致（missing-surface）：文档把uniform列为 Uniform 地址空间的构造，但只给出了ConstantBuffer<T>的常见用法示例。实际上编译器把入口点参数上的uniform映射为 SPIR-V 的PushConstant存储类，而非ConstantBuffer<T>的Uniform存储类。测试束建议文档补充说明：Vulkan/SPIR-V 上入口点参数的uniform可能降级为 push-constant 块，而ConstantBuffer<T>恒定映射到Uniform存储类。对应测试 uniform-keyword-spirv.slang。

2. static全局变量在 SPIR-V 上的存储类未文档化（missing-surface）：编译器在变量存活时将其发射为Private（逐调用实例），但一旦所有使用点内联进同一个函数，就可能被提升为 SSA 值或Function作用域存储——可观察的存储类取决于内联情况而非声明本身。对应测试 thread-local-static-hlsl.slang 中的SPIRV: %g_perThread = OpVariable {{.*}} Private断言。

3. Remark 2 的"generally"存在歧义（ambiguous-claim）：文档说跨地址空间指针"通常"不可互换，但未说明混用地址空间到底是编译期错误、类型系统错误还是未定义运行时行为；"通常"一词也无法回答是否存在合法的跨空间指针转换（例如通过reinterpret_cast或Ptr<T>）。测试束建议文档明确：跨地址空间指针转换是否恒为类型错误（若是，给出 E#### 诊断码）、恒为 UB，还是在特定文档化场景下被允许。

这三条观察体现了"以测试验证文档、以文档反哺测试"的闭环：每个缺口都锚定到 docs/language-reference/basics-memory-model-address-spaces.md#address-spaces 小节，并给出可落地的文档修改建议。

七、实践要点总结

• 选空间即选性能与语义：Uniform/Push constant/Storage buffer 是跨线程共享数据的主流通道；Group-shared 是线程组内通信的显存级高速通道（配合屏障使用）；Thread-local 与 Function 空间适合线程私有状态；Input/Output 用于图形管线阶段间数据；Specialization constant 用于管线创建期定值。
• 同一份 Slang 代码可以无改动地发射到 6 种后端，但各后端的存储类名称完全不同（如Workgroup/threadgroup/__shared__/shared/groupshared/var<workgroup>），这正体现了 Slang 地址空间抽象的价值：开发者面向 Slang 的 11 种地址空间编程，由编译器负责与 SPIR-V 存储类、Metal 地址空间限定符、WGSL address space、HLSL 寄存器语义逐一对齐。
• 验证手段：任何地址空间相关的改动都可以用本束的模式自测——编写一段最小着色器，用//TEST:SIMPLE(filecheck=...)配合spirv-asm/metal/wgsl/glsl/hlsl（必要时加cuda）多目标发射，再用CHECK/CHECK-DAG钉定目标存储类；对 GPU 专属语义（如跨线程组指针别名）则需在带独立 GPU 的 CI 上验证。

如需进一步了解与地址空间紧密相关的运行语义（如内存模型一致性、内存屏障），可继续阅读语言参考中的 basics-memory-model.md、basics-memory-model-consistency.md 与 basics-memory-model-special-topics.md，以及与图形管线阶段相关的 shaders-and-kernels.md。

【免费下载链接】slangMaking it easier to work with shaders项目地址: https://gitcode.com/GitHub_Trending/sl/slang