From 0ce98d6411ddadd078aaf3b5b1fff0b17be3b867 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Fri, 4 Feb 2022 13:31:16 -0800 Subject: [PATCH] [msl-out][spv-out][glsl-out][hlsl-out] Fix ArraySize on globals. --- src/back/glsl/mod.rs | 6 + src/back/hlsl/help.rs | 1 + src/back/msl/writer.rs | 1 + src/back/spv/index.rs | 23 +++- src/back/spv/mod.rs | 15 ++- src/lib.rs | 3 +- tests/in/globals.wgsl | 1 + tests/out/glsl/globals.main.Compute.glsl | 3 + tests/out/hlsl/globals.hlsl | 8 ++ tests/out/msl/globals.msl | 3 + tests/out/spv/globals.spvasm | 158 ++++++++++++----------- tests/out/wgsl/globals.wgsl | 1 + 12 files changed, 140 insertions(+), 83 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 38888957eb..20d0035037 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -129,6 +129,12 @@ impl<'a> GlobalTypeKind<'a> { } => Self::Unsized(members), _ => Self::WrappedStruct, }, + // Naga IR permits globals to be dynamically sized arrays. Render + // these in GLSL as buffers. + crate::TypeInner::Array { + size: crate::ArraySize::Dynamic, + .. + } => Self::WrappedStruct, _ => Self::Other, } } diff --git a/src/back/hlsl/help.rs b/src/back/hlsl/help.rs index 8c4b31e802..428df8afe5 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -432,6 +432,7 @@ impl<'a, W: Write> super::Writer<'a, W> { match func_ctx.expressions[handle] { crate::Expression::ArrayLength(expr) => { let global_expr = match func_ctx.expressions[expr] { + crate::Expression::GlobalVariable(_) => expr, crate::Expression::AccessIndex { base, index: _ } => base, ref other => unreachable!("Array length of {:?}", other), }; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 7da11e6bee..f341a38ff2 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1397,6 +1397,7 @@ impl Writer { _ => return Err(Error::Validation), } } + crate::Expression::GlobalVariable(handle) => handle, _ => return Err(Error::Validation), }; diff --git a/src/back/spv/index.rs b/src/back/spv/index.rs index fa36d25289..d2cbdf4d6d 100644 --- a/src/back/spv/index.rs +++ b/src/back/spv/index.rs @@ -2,7 +2,10 @@ Bounds-checking for SPIR-V output. */ -use super::{selection::Selection, Block, BlockContext, Error, IdGenerator, Instruction, Word}; +use super::{ + helpers::global_needs_wrapper, selection::Selection, Block, BlockContext, Error, IdGenerator, + Instruction, Word, +}; use crate::{arena::Handle, proc::BoundsCheckPolicy}; /// The results of performing a bounds check. @@ -32,16 +35,18 @@ pub(super) enum MaybeKnown { impl<'w> BlockContext<'w> { /// Emit code to compute the length of a run-time array. /// - /// Given `array`, an expression referring to the final member of a struct, - /// where the member in question is a runtime-sized array, return the + /// Given `array`, an expression referring a runtime-sized array, return the /// instruction id for the array's length. pub(super) fn write_runtime_array_length( &mut self, array: Handle, block: &mut Block, ) -> Result { - // Look into the expression to find the value and type of the struct - // holding the dynamically-sized array. + // Naga IR permits runtime-sized arrays as global variables or as the + // final member of a struct that is a global variable. SPIR-V permits + // only the latter, so this back end wraps bare runtime-sized arrays + // in a made-up struct; see `helpers::global_needs_wrapper` and its uses. + // This code must handle both cases. let (structure_id, last_member_index) = match self.ir_function.expressions[array] { crate::Expression::AccessIndex { base, index } => { match self.ir_function.expressions[base] { @@ -52,6 +57,14 @@ impl<'w> BlockContext<'w> { _ => return Err(Error::Validation("array length expression")), } } + crate::Expression::GlobalVariable(handle) => { + let global = &self.ir_module.global_variables[handle]; + if !global_needs_wrapper(self.ir_module, global) { + return Err(Error::Validation("array length expression")); + } + + (self.writer.global_variables[handle.index()].var_id, 0) + } _ => return Err(Error::Validation("array length expression")), }; diff --git a/src/back/spv/mod.rs b/src/back/spv/mod.rs index 0d03527d68..538147ed16 100644 --- a/src/back/spv/mod.rs +++ b/src/back/spv/mod.rs @@ -433,12 +433,25 @@ impl recyclable::Recyclable for CachedExpressions { #[derive(Clone)] struct GlobalVariable { - /// ID of the variable. Not really used. + /// ID of the OpVariable that declares the global. + /// + /// If you need the variable's value, use [`access_id`] instead of this + /// field. If we wrapped the Naga IR `GlobalVariable`'s type in a struct to + /// comply with Vulkan's requirements, then this points to the `OpVariable` + /// with the synthesized struct type, whereas `access_id` points to the + /// field of said struct that holds the variable's actual value. + /// + /// This is used to compute the `access_id` pointer in function prologues, + /// and used for `ArrayLength` expressions, which do need the struct. + /// + /// [`access_id`]: GlobalVariable::access_id var_id: Word, + /// For `AddressSpace::Handle` variables, this ID is recorded in the function /// prelude block (and reset before every function) as `OpLoad` of the variable. /// It is then used for all the global ops, such as `OpImageSample`. handle_id: Word, + /// Actual ID used to access this variable. /// For wrapped buffer variables, this ID is `OpAccessChain` into the /// wrapper. Otherwise, the same as `var_id`. diff --git a/src/lib.rs b/src/lib.rs index f384c6ea14..ac1118c060 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -652,7 +652,8 @@ pub enum TypeInner { /// An `Array` is [`SIZED`] unless its `size` is [`Dynamic`]. /// Dynamically-sized arrays may only appear in a few situations: /// - /// - They may appear as the last member of a [`Struct`]. + /// - They may appear as the type of a [`GlobalVariable`], or as the last + /// member of a [`Struct`]. /// /// - They may appear as the base type of a [`Pointer`]. An /// [`AccessIndex`] expression referring to a struct's final diff --git a/tests/in/globals.wgsl b/tests/in/globals.wgsl index 23dabf1b1b..e9f3fa77eb 100644 --- a/tests/in/globals.wgsl +++ b/tests/in/globals.wgsl @@ -23,6 +23,7 @@ var float_vecs: array, 20>; fn main() { wg[3] = alignment.v1; wg[2] = alignment.v3.x; + wg[1] = f32(arrayLength(&dummy)); atomicStore(&at, 2u); // Valid, Foo and at is in function scope diff --git a/tests/out/glsl/globals.main.Compute.glsl b/tests/out/glsl/globals.main.Compute.glsl index a40e6bbac9..46150907be 100644 --- a/tests/out/glsl/globals.main.Compute.glsl +++ b/tests/out/glsl/globals.main.Compute.glsl @@ -15,6 +15,8 @@ shared uint at_1; layout(std430) readonly buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; }; +layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; }; + void main() { float Foo_1 = 1.0; @@ -23,6 +25,7 @@ void main() { wg[3] = _e9; float _e14 = _group_0_binding_1_cs.v3_.x; wg[2] = _e14; + wg[1] = float(uint(_group_0_binding_2_cs.length())); at_1 = 2u; return; } diff --git a/tests/out/hlsl/globals.hlsl b/tests/out/hlsl/globals.hlsl index a6b91c9523..515ff1de6b 100644 --- a/tests/out/hlsl/globals.hlsl +++ b/tests/out/hlsl/globals.hlsl @@ -11,6 +11,13 @@ ByteAddressBuffer alignment : register(t1); ByteAddressBuffer dummy : register(t2); cbuffer float_vecs : register(b3) { float4 float_vecs[20]; } +uint NagaBufferLength(ByteAddressBuffer buffer) +{ + uint ret; + buffer.GetDimensions(ret); + return ret; +} + [numthreads(1, 1, 1)] void main() { @@ -21,6 +28,7 @@ void main() wg[3] = _expr9; float _expr14 = asfloat(alignment.Load(0+0)); wg[2] = _expr14; + wg[1] = float(((NagaBufferLength(dummy) - 0) / 8)); at_1 = 2u; return; } diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index c8c58b3391..b56bc647dc 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -23,6 +23,8 @@ kernel void main_( threadgroup type_2& wg , threadgroup metal::atomic_uint& at_1 , device Foo& alignment [[user(fake0)]] +, device type_6& dummy [[user(fake0)]] +, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { float Foo_1 = 1.0; bool at = true; @@ -30,6 +32,7 @@ kernel void main_( wg.inner[3] = _e9; float _e14 = metal::float3(alignment.v3_).x; wg.inner[2] = _e14; + wg.inner[1] = static_cast(1 + (_buffer_sizes.size3 - 0 - 8) / 8); metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed); return; } diff --git a/tests/out/spv/globals.spvasm b/tests/out/spv/globals.spvasm index ecc2805ef4..bfba6116fa 100644 --- a/tests/out/spv/globals.spvasm +++ b/tests/out/spv/globals.spvasm @@ -1,32 +1,32 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 61 +; Bound: 66 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %40 "main" -OpExecutionMode %40 LocalSize 1 1 1 -OpDecorate %15 ArrayStride 4 -OpMemberDecorate %17 0 Offset 0 -OpMemberDecorate %17 1 Offset 12 -OpDecorate %19 ArrayStride 8 -OpDecorate %21 ArrayStride 16 -OpDecorate %26 NonWritable -OpDecorate %26 DescriptorSet 0 -OpDecorate %26 Binding 1 -OpDecorate %27 Block -OpMemberDecorate %27 0 Offset 0 -OpDecorate %29 NonWritable -OpDecorate %29 DescriptorSet 0 -OpDecorate %29 Binding 2 -OpDecorate %30 Block -OpMemberDecorate %30 0 Offset 0 -OpDecorate %32 DescriptorSet 0 -OpDecorate %32 Binding 3 -OpDecorate %33 Block -OpMemberDecorate %33 0 Offset 0 +OpEntryPoint GLCompute %41 "main" +OpExecutionMode %41 LocalSize 1 1 1 +OpDecorate %16 ArrayStride 4 +OpMemberDecorate %18 0 Offset 0 +OpMemberDecorate %18 1 Offset 12 +OpDecorate %20 ArrayStride 8 +OpDecorate %22 ArrayStride 16 +OpDecorate %27 NonWritable +OpDecorate %27 DescriptorSet 0 +OpDecorate %27 Binding 1 +OpDecorate %28 Block +OpMemberDecorate %28 0 Offset 0 +OpDecorate %30 NonWritable +OpDecorate %30 DescriptorSet 0 +OpDecorate %30 Binding 2 +OpDecorate %31 Block +OpMemberDecorate %31 0 Offset 0 +OpDecorate %33 DescriptorSet 0 +OpDecorate %33 Binding 3 +OpDecorate %34 Block +OpMemberDecorate %34 0 Offset 0 %2 = OpTypeVoid %4 = OpTypeBool %3 = OpConstantTrue %4 @@ -36,59 +36,65 @@ OpMemberDecorate %33 0 Offset 0 %7 = OpConstant %8 20 %9 = OpConstant %8 3 %10 = OpConstant %8 2 -%11 = OpConstant %6 2 -%13 = OpTypeFloat 32 -%12 = OpConstant %13 1.0 -%14 = OpConstantTrue %4 -%15 = OpTypeArray %13 %5 -%16 = OpTypeVector %13 3 -%17 = OpTypeStruct %16 %13 -%18 = OpTypeVector %13 2 -%19 = OpTypeRuntimeArray %18 -%20 = OpTypeVector %13 4 -%21 = OpTypeArray %20 %7 -%23 = OpTypePointer Workgroup %15 -%22 = OpVariable %23 Workgroup -%25 = OpTypePointer Workgroup %6 -%24 = OpVariable %25 Workgroup -%27 = OpTypeStruct %17 -%28 = OpTypePointer StorageBuffer %27 -%26 = OpVariable %28 StorageBuffer -%30 = OpTypeStruct %19 -%31 = OpTypePointer StorageBuffer %30 -%29 = OpVariable %31 StorageBuffer -%33 = OpTypeStruct %21 -%34 = OpTypePointer Uniform %33 -%32 = OpVariable %34 Uniform -%36 = OpTypePointer Function %13 -%38 = OpTypePointer Function %4 -%41 = OpTypeFunction %2 -%42 = OpTypePointer StorageBuffer %17 -%43 = OpConstant %6 0 -%45 = OpTypePointer StorageBuffer %19 -%46 = OpTypePointer Uniform %21 -%48 = OpTypePointer Workgroup %13 -%49 = OpTypePointer StorageBuffer %13 -%50 = OpConstant %6 1 -%53 = OpConstant %6 3 -%55 = OpTypePointer StorageBuffer %16 -%56 = OpTypePointer StorageBuffer %13 -%60 = OpConstant %6 256 -%40 = OpFunction %2 None %41 -%39 = OpLabel -%35 = OpVariable %36 Function %12 -%37 = OpVariable %38 Function %14 -%44 = OpAccessChain %42 %26 %43 -OpBranch %47 -%47 = OpLabel -%51 = OpAccessChain %49 %44 %50 -%52 = OpLoad %13 %51 -%54 = OpAccessChain %48 %22 %53 -OpStore %54 %52 -%57 = OpAccessChain %56 %44 %43 %43 -%58 = OpLoad %13 %57 -%59 = OpAccessChain %48 %22 %11 -OpStore %59 %58 -OpAtomicStore %24 %10 %60 %11 +%11 = OpConstant %8 1 +%12 = OpConstant %6 2 +%14 = OpTypeFloat 32 +%13 = OpConstant %14 1.0 +%15 = OpConstantTrue %4 +%16 = OpTypeArray %14 %5 +%17 = OpTypeVector %14 3 +%18 = OpTypeStruct %17 %14 +%19 = OpTypeVector %14 2 +%20 = OpTypeRuntimeArray %19 +%21 = OpTypeVector %14 4 +%22 = OpTypeArray %21 %7 +%24 = OpTypePointer Workgroup %16 +%23 = OpVariable %24 Workgroup +%26 = OpTypePointer Workgroup %6 +%25 = OpVariable %26 Workgroup +%28 = OpTypeStruct %18 +%29 = OpTypePointer StorageBuffer %28 +%27 = OpVariable %29 StorageBuffer +%31 = OpTypeStruct %20 +%32 = OpTypePointer StorageBuffer %31 +%30 = OpVariable %32 StorageBuffer +%34 = OpTypeStruct %22 +%35 = OpTypePointer Uniform %34 +%33 = OpVariable %35 Uniform +%37 = OpTypePointer Function %14 +%39 = OpTypePointer Function %4 +%42 = OpTypeFunction %2 +%43 = OpTypePointer StorageBuffer %18 +%44 = OpConstant %6 0 +%46 = OpTypePointer StorageBuffer %20 +%48 = OpTypePointer Uniform %22 +%50 = OpTypePointer Workgroup %14 +%51 = OpTypePointer StorageBuffer %14 +%52 = OpConstant %6 1 +%55 = OpConstant %6 3 +%57 = OpTypePointer StorageBuffer %17 +%58 = OpTypePointer StorageBuffer %14 +%65 = OpConstant %6 256 +%41 = OpFunction %2 None %42 +%40 = OpLabel +%36 = OpVariable %37 Function %13 +%38 = OpVariable %39 Function %15 +%45 = OpAccessChain %43 %27 %44 +%47 = OpAccessChain %46 %30 %44 +OpBranch %49 +%49 = OpLabel +%53 = OpAccessChain %51 %45 %52 +%54 = OpLoad %14 %53 +%56 = OpAccessChain %50 %23 %55 +OpStore %56 %54 +%59 = OpAccessChain %58 %45 %44 %44 +%60 = OpLoad %14 %59 +%61 = OpAccessChain %50 %23 %12 +OpStore %61 %60 +%62 = OpArrayLength %6 %30 0 +%63 = OpConvertUToF %14 %62 +%64 = OpAccessChain %50 %23 %52 +OpStore %64 %63 +OpAtomicStore %25 %10 %65 %12 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/globals.wgsl b/tests/out/wgsl/globals.wgsl index 3196468c50..f73c6c04d3 100644 --- a/tests/out/wgsl/globals.wgsl +++ b/tests/out/wgsl/globals.wgsl @@ -23,6 +23,7 @@ fn main() { wg[3] = _e9; let _e14 = alignment.v3_.x; wg[2] = _e14; + wg[1] = f32(arrayLength((&dummy))); atomicStore((&at_1), 2u); return; }