From 27d0fee2e7cd4b7e21b05f083f5d0277acddd98c Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 12 Apr 2022 18:54:42 -0700 Subject: [PATCH] More tests for non-struct global types. --- src/back/glsl/mod.rs | 6 +- tests/in/globals.wgsl | 6 + tests/out/glsl/globals.main.Compute.glsl | 22 +- tests/out/hlsl/globals.hlsl | 19 +- tests/out/msl/globals.msl | 20 +- tests/out/spv/globals.spvasm | 262 +++++++++++++---------- tests/out/wgsl/globals.wgsl | 20 +- 7 files changed, 214 insertions(+), 141 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 47cbd920f..c4d66afd6 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -923,13 +923,17 @@ impl<'a, W: Write> Writer<'a, W> { crate::AddressSpace::PushConstant => { self.write_simple_global(handle, global)?; } - crate::AddressSpace::Uniform | crate::AddressSpace::Handle => { + crate::AddressSpace::Uniform => { self.write_interface_block(handle, global)?; } crate::AddressSpace::Storage { .. } => { self.write_interface_block(handle, global)?; } + // A global variable in the `Function` address space is a + // contradiction in terms. crate::AddressSpace::Function => unreachable!(), + // Textures and samplers are handled directly in `Writer::write`. + crate::AddressSpace::Handle => unreachable!(), } Ok(()) diff --git a/tests/in/globals.wgsl b/tests/in/globals.wgsl index 86a41569d..573c1575c 100644 --- a/tests/in/globals.wgsl +++ b/tests/in/globals.wgsl @@ -19,6 +19,9 @@ var dummy: array>; @group(0) @binding(3) var float_vecs: array, 20>; +@group(0) @binding(4) +var global_vec: vec4; + fn test_msl_packed_vec3_as_arg(arg: vec3) {} fn test_msl_packed_vec3() { @@ -50,6 +53,9 @@ fn test_msl_packed_vec3() { fn main() { test_msl_packed_vec3(); + wg[6] = global_vec.x; + wg[5] = dummy[1].y; + wg[4] = float_vecs[0].w; wg[3] = alignment.v1; wg[2] = alignment.v3.x; alignment.v1 = 4.0; diff --git a/tests/out/glsl/globals.main.Compute.glsl b/tests/out/glsl/globals.main.Compute.glsl index 0db8ee14b..785a3d1e1 100644 --- a/tests/out/glsl/globals.main.Compute.glsl +++ b/tests/out/glsl/globals.main.Compute.glsl @@ -17,6 +17,10 @@ layout(std430) buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; }; layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; }; +uniform type_8_block_2Compute { vec4 _group_0_binding_3_cs[20]; }; + +uniform type_7_block_3Compute { vec4 _group_0_binding_4_cs; }; + void test_msl_packed_vec3_as_arg(vec3 arg) { return; @@ -27,8 +31,8 @@ void test_msl_packed_vec3_() { _group_0_binding_1_cs.v3_ = vec3(1.0); _group_0_binding_1_cs.v3_.x = 1.0; _group_0_binding_1_cs.v3_.x = 2.0; - int _e19 = idx; - _group_0_binding_1_cs.v3_[_e19] = 3.0; + int _e20 = idx; + _group_0_binding_1_cs.v3_[_e20] = 3.0; Foo data = _group_0_binding_1_cs; vec3 unnamed = data.v3_; vec2 unnamed_1 = data.v3_.zx; @@ -43,10 +47,16 @@ void main() { float Foo_1 = 1.0; bool at = true; test_msl_packed_vec3_(); - float _e9 = _group_0_binding_1_cs.v1_; - wg[3] = _e9; - float _e14 = _group_0_binding_1_cs.v3_.x; - wg[2] = _e14; + float _e10 = _group_0_binding_4_cs.x; + wg[6] = _e10; + float _e16 = _group_0_binding_2_cs[1].y; + wg[5] = _e16; + float _e22 = _group_0_binding_3_cs[0].w; + wg[4] = _e22; + float _e26 = _group_0_binding_1_cs.v1_; + wg[3] = _e26; + float _e31 = _group_0_binding_1_cs.v3_.x; + wg[2] = _e31; _group_0_binding_1_cs.v1_ = 4.0; wg[1] = float(uint(_group_0_binding_2_cs.length())); at_1 = 2u; diff --git a/tests/out/hlsl/globals.hlsl b/tests/out/hlsl/globals.hlsl index 7844ddf35..abf543ff6 100644 --- a/tests/out/hlsl/globals.hlsl +++ b/tests/out/hlsl/globals.hlsl @@ -10,6 +10,7 @@ groupshared uint at_1; RWByteAddressBuffer alignment : register(u1); ByteAddressBuffer dummy : register(t2); cbuffer float_vecs : register(b3) { float4 float_vecs[20]; } +cbuffer global_vec : register(b4) { float4 global_vec; } void test_msl_packed_vec3_as_arg(float3 arg) { @@ -23,8 +24,8 @@ void test_msl_packed_vec3_() alignment.Store3(0, asuint(float3(1.0.xxx))); alignment.Store(0+0, asuint(1.0)); alignment.Store(0+0, asuint(2.0)); - int _expr19 = idx; - alignment.Store(_expr19*4+0, asuint(3.0)); + int _expr20 = idx; + alignment.Store(_expr20*4+0, asuint(3.0)); Foo data = {asfloat(alignment.Load3(0)), asfloat(alignment.Load(12))}; float3 unnamed = data.v3_; float2 unnamed_1 = data.v3_.zx; @@ -49,10 +50,16 @@ void main() bool at = true; test_msl_packed_vec3_(); - float _expr9 = asfloat(alignment.Load(12)); - wg[3] = _expr9; - float _expr14 = asfloat(alignment.Load(0+0)); - wg[2] = _expr14; + float _expr10 = global_vec.x; + wg[6] = _expr10; + float _expr16 = asfloat(dummy.Load(4+8)); + wg[5] = _expr16; + float _expr22 = float_vecs[0].w; + wg[4] = _expr22; + float _expr26 = asfloat(alignment.Load(12)); + wg[3] = _expr26; + float _expr31 = asfloat(alignment.Load(0+0)); + wg[2] = _expr31; alignment.Store(12, asuint(4.0)); wg[1] = float(((NagaBufferLength(dummy) - 0) / 8)); at_1 = 2u; diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index 504e41d57..3c83f67dd 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -36,8 +36,8 @@ void test_msl_packed_vec3_( alignment.v3_ = metal::float3(1.0); alignment.v3_[0] = 1.0; alignment.v3_[0] = 2.0; - int _e19 = idx; - alignment.v3_[_e19] = 3.0; + int _e20 = idx; + alignment.v3_[_e20] = 3.0; Foo data = alignment; metal::float3 unnamed = data.v3_; metal::float2 unnamed_1 = metal::float3(data.v3_).zx; @@ -53,15 +53,23 @@ kernel void main_( , threadgroup metal::atomic_uint& at_1 , device Foo& alignment [[user(fake0)]] , device type_6 const& dummy [[user(fake0)]] +, constant type_8& float_vecs [[user(fake0)]] +, constant metal::float4& global_vec [[user(fake0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { float Foo_1 = 1.0; bool at = true; test_msl_packed_vec3_(alignment); - float _e9 = alignment.v1_; - wg.inner[3] = _e9; - float _e14 = alignment.v3_[0]; - wg.inner[2] = _e14; + float _e10 = global_vec.x; + wg.inner[6] = _e10; + float _e16 = dummy[1].y; + wg.inner[5] = _e16; + float _e22 = float_vecs.inner[0].w; + wg.inner[4] = _e22; + float _e26 = alignment.v1_; + wg.inner[3] = _e26; + float _e31 = alignment.v3_[0]; + wg.inner[2] = _e31; alignment.v1_ = 4.0; wg.inner[1] = static_cast(1 + (_buffer_sizes.size3 - 0 - 8) / 8); metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed); diff --git a/tests/out/spv/globals.spvasm b/tests/out/spv/globals.spvasm index 395096db3..07e6da2a7 100644 --- a/tests/out/spv/globals.spvasm +++ b/tests/out/spv/globals.spvasm @@ -1,31 +1,35 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 107 +; Bound: 130 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %87 "main" -OpExecutionMode %87 LocalSize 1 1 1 -OpDecorate %21 ArrayStride 4 -OpMemberDecorate %23 0 Offset 0 -OpMemberDecorate %23 1 Offset 12 -OpDecorate %25 ArrayStride 8 -OpDecorate %27 ArrayStride 16 -OpDecorate %35 DescriptorSet 0 -OpDecorate %35 Binding 1 -OpDecorate %36 Block -OpMemberDecorate %36 0 Offset 0 -OpDecorate %38 NonWritable +OpEntryPoint GLCompute %94 "main" +OpExecutionMode %94 LocalSize 1 1 1 +OpDecorate %24 ArrayStride 4 +OpMemberDecorate %26 0 Offset 0 +OpMemberDecorate %26 1 Offset 12 +OpDecorate %28 ArrayStride 8 +OpDecorate %30 ArrayStride 16 OpDecorate %38 DescriptorSet 0 -OpDecorate %38 Binding 2 +OpDecorate %38 Binding 1 OpDecorate %39 Block OpMemberDecorate %39 0 Offset 0 +OpDecorate %41 NonWritable OpDecorate %41 DescriptorSet 0 -OpDecorate %41 Binding 3 +OpDecorate %41 Binding 2 OpDecorate %42 Block OpMemberDecorate %42 0 Offset 0 +OpDecorate %44 DescriptorSet 0 +OpDecorate %44 Binding 3 +OpDecorate %45 Block +OpMemberDecorate %45 0 Offset 0 +OpDecorate %47 DescriptorSet 0 +OpDecorate %47 Binding 4 +OpDecorate %48 Block +OpMemberDecorate %48 0 Offset 0 %2 = OpTypeVoid %4 = OpTypeBool %3 = OpConstantTrue %4 @@ -40,112 +44,138 @@ OpMemberDecorate %42 0 Offset 0 %13 = OpConstant %10 2.0 %14 = OpConstant %10 3.0 %15 = OpConstant %10 0.0 -%16 = OpConstant %8 3 -%17 = OpConstant %8 2 -%18 = OpConstant %10 4.0 -%19 = OpConstant %6 2 -%20 = OpConstantTrue %4 -%21 = OpTypeArray %10 %5 -%22 = OpTypeVector %10 3 -%23 = OpTypeStruct %22 %10 -%24 = OpTypeVector %10 2 -%25 = OpTypeRuntimeArray %24 -%26 = OpTypeVector %10 4 -%27 = OpTypeArray %26 %7 -%28 = OpTypeMatrix %22 3 -%29 = OpConstantComposite %22 %15 %15 %15 -%30 = OpConstantComposite %28 %29 %29 %29 -%32 = OpTypePointer Workgroup %21 -%31 = OpVariable %32 Workgroup -%34 = OpTypePointer Workgroup %6 -%33 = OpVariable %34 Workgroup -%36 = OpTypeStruct %23 -%37 = OpTypePointer StorageBuffer %36 -%35 = OpVariable %37 StorageBuffer -%39 = OpTypeStruct %25 +%16 = OpConstant %8 6 +%17 = OpConstant %8 5 +%18 = OpConstant %8 4 +%19 = OpConstant %8 3 +%20 = OpConstant %8 2 +%21 = OpConstant %10 4.0 +%22 = OpConstant %6 2 +%23 = OpConstantTrue %4 +%24 = OpTypeArray %10 %5 +%25 = OpTypeVector %10 3 +%26 = OpTypeStruct %25 %10 +%27 = OpTypeVector %10 2 +%28 = OpTypeRuntimeArray %27 +%29 = OpTypeVector %10 4 +%30 = OpTypeArray %29 %7 +%31 = OpTypeMatrix %25 3 +%32 = OpConstantComposite %25 %15 %15 %15 +%33 = OpConstantComposite %31 %32 %32 %32 +%35 = OpTypePointer Workgroup %24 +%34 = OpVariable %35 Workgroup +%37 = OpTypePointer Workgroup %6 +%36 = OpVariable %37 Workgroup +%39 = OpTypeStruct %26 %40 = OpTypePointer StorageBuffer %39 %38 = OpVariable %40 StorageBuffer -%42 = OpTypeStruct %27 -%43 = OpTypePointer Uniform %42 -%41 = OpVariable %43 Uniform -%47 = OpTypeFunction %2 %22 -%48 = OpTypePointer StorageBuffer %23 -%49 = OpTypePointer StorageBuffer %25 -%50 = OpTypePointer Uniform %27 -%53 = OpTypePointer Function %8 -%56 = OpTypeFunction %2 -%57 = OpConstant %6 0 -%60 = OpTypePointer StorageBuffer %22 -%63 = OpTypePointer StorageBuffer %10 -%83 = OpTypePointer Function %10 -%85 = OpTypePointer Function %4 -%92 = OpTypePointer Workgroup %10 -%93 = OpTypePointer StorageBuffer %10 -%94 = OpConstant %6 1 -%97 = OpConstant %6 3 -%106 = OpConstant %6 256 -%46 = OpFunction %2 None %47 -%45 = OpFunctionParameter %22 -%44 = OpLabel -OpBranch %51 -%51 = OpLabel +%42 = OpTypeStruct %28 +%43 = OpTypePointer StorageBuffer %42 +%41 = OpVariable %43 StorageBuffer +%45 = OpTypeStruct %30 +%46 = OpTypePointer Uniform %45 +%44 = OpVariable %46 Uniform +%48 = OpTypeStruct %29 +%49 = OpTypePointer Uniform %48 +%47 = OpVariable %49 Uniform +%53 = OpTypeFunction %2 %25 +%54 = OpTypePointer StorageBuffer %26 +%55 = OpTypePointer StorageBuffer %28 +%56 = OpTypePointer Uniform %29 +%57 = OpTypePointer Uniform %30 +%60 = OpTypePointer Function %8 +%63 = OpTypeFunction %2 +%64 = OpConstant %6 0 +%67 = OpTypePointer StorageBuffer %25 +%70 = OpTypePointer StorageBuffer %10 +%90 = OpTypePointer Function %10 +%92 = OpTypePointer Function %4 +%101 = OpTypePointer Workgroup %10 +%102 = OpTypePointer Uniform %10 +%105 = OpConstant %6 6 +%107 = OpTypePointer StorageBuffer %27 +%108 = OpConstant %6 1 +%111 = OpConstant %6 5 +%113 = OpConstant %6 3 +%116 = OpConstant %6 4 +%118 = OpTypePointer StorageBuffer %10 +%129 = OpConstant %6 256 +%52 = OpFunction %2 None %53 +%51 = OpFunctionParameter %25 +%50 = OpLabel +OpBranch %58 +%58 = OpLabel OpReturn OpFunctionEnd -%55 = OpFunction %2 None %56 -%54 = OpLabel -%52 = OpVariable %53 Function %11 -%58 = OpAccessChain %48 %35 %57 -OpBranch %59 -%59 = OpLabel -%61 = OpCompositeConstruct %22 %9 %9 %9 -%62 = OpAccessChain %60 %58 %57 -OpStore %62 %61 -%64 = OpAccessChain %63 %58 %57 %57 -OpStore %64 %9 -%65 = OpAccessChain %63 %58 %57 %57 -OpStore %65 %13 -%66 = OpLoad %8 %52 -%67 = OpAccessChain %63 %58 %57 %66 -OpStore %67 %14 -%68 = OpLoad %23 %58 -%69 = OpCompositeExtract %22 %68 0 -%70 = OpCompositeExtract %22 %68 0 -%71 = OpVectorShuffle %24 %70 %70 2 0 -%72 = OpCompositeExtract %22 %68 0 -%73 = OpFunctionCall %2 %46 %72 -%74 = OpCompositeExtract %22 %68 0 -%75 = OpVectorTimesMatrix %22 %74 %30 -%76 = OpCompositeExtract %22 %68 0 -%77 = OpMatrixTimesVector %22 %30 %76 -%78 = OpCompositeExtract %22 %68 0 -%79 = OpVectorTimesScalar %22 %78 %13 -%80 = OpCompositeExtract %22 %68 0 -%81 = OpVectorTimesScalar %22 %80 %13 +%62 = OpFunction %2 None %63 +%61 = OpLabel +%59 = OpVariable %60 Function %11 +%65 = OpAccessChain %54 %38 %64 +OpBranch %66 +%66 = OpLabel +%68 = OpCompositeConstruct %25 %9 %9 %9 +%69 = OpAccessChain %67 %65 %64 +OpStore %69 %68 +%71 = OpAccessChain %70 %65 %64 %64 +OpStore %71 %9 +%72 = OpAccessChain %70 %65 %64 %64 +OpStore %72 %13 +%73 = OpLoad %8 %59 +%74 = OpAccessChain %70 %65 %64 %73 +OpStore %74 %14 +%75 = OpLoad %26 %65 +%76 = OpCompositeExtract %25 %75 0 +%77 = OpCompositeExtract %25 %75 0 +%78 = OpVectorShuffle %27 %77 %77 2 0 +%79 = OpCompositeExtract %25 %75 0 +%80 = OpFunctionCall %2 %52 %79 +%81 = OpCompositeExtract %25 %75 0 +%82 = OpVectorTimesMatrix %25 %81 %33 +%83 = OpCompositeExtract %25 %75 0 +%84 = OpMatrixTimesVector %25 %33 %83 +%85 = OpCompositeExtract %25 %75 0 +%86 = OpVectorTimesScalar %25 %85 %13 +%87 = OpCompositeExtract %25 %75 0 +%88 = OpVectorTimesScalar %25 %87 %13 OpReturn OpFunctionEnd -%87 = OpFunction %2 None %56 -%86 = OpLabel -%82 = OpVariable %83 Function %9 -%84 = OpVariable %85 Function %20 -%88 = OpAccessChain %48 %35 %57 -%89 = OpAccessChain %49 %38 %57 -OpBranch %90 -%90 = OpLabel -%91 = OpFunctionCall %2 %55 -%95 = OpAccessChain %93 %88 %94 -%96 = OpLoad %10 %95 -%98 = OpAccessChain %92 %31 %97 -OpStore %98 %96 -%99 = OpAccessChain %63 %88 %57 %57 -%100 = OpLoad %10 %99 -%101 = OpAccessChain %92 %31 %19 -OpStore %101 %100 -%102 = OpAccessChain %93 %88 %94 -OpStore %102 %18 -%103 = OpArrayLength %6 %38 0 -%104 = OpConvertUToF %10 %103 -%105 = OpAccessChain %92 %31 %94 -OpStore %105 %104 -OpAtomicStore %33 %17 %106 %19 +%94 = OpFunction %2 None %63 +%93 = OpLabel +%89 = OpVariable %90 Function %9 +%91 = OpVariable %92 Function %23 +%95 = OpAccessChain %54 %38 %64 +%96 = OpAccessChain %55 %41 %64 +%97 = OpAccessChain %57 %44 %64 +%98 = OpAccessChain %56 %47 %64 +OpBranch %99 +%99 = OpLabel +%100 = OpFunctionCall %2 %62 +%103 = OpAccessChain %102 %98 %64 +%104 = OpLoad %10 %103 +%106 = OpAccessChain %101 %34 %105 +OpStore %106 %104 +%109 = OpAccessChain %70 %96 %108 %108 +%110 = OpLoad %10 %109 +%112 = OpAccessChain %101 %34 %111 +OpStore %112 %110 +%114 = OpAccessChain %102 %97 %64 %113 +%115 = OpLoad %10 %114 +%117 = OpAccessChain %101 %34 %116 +OpStore %117 %115 +%119 = OpAccessChain %118 %95 %108 +%120 = OpLoad %10 %119 +%121 = OpAccessChain %101 %34 %113 +OpStore %121 %120 +%122 = OpAccessChain %70 %95 %64 %64 +%123 = OpLoad %10 %122 +%124 = OpAccessChain %101 %34 %22 +OpStore %124 %123 +%125 = OpAccessChain %118 %95 %108 +OpStore %125 %21 +%126 = OpArrayLength %6 %41 0 +%127 = OpConvertUToF %10 %126 +%128 = OpAccessChain %101 %34 %108 +OpStore %128 %127 +OpAtomicStore %36 %20 %129 %22 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/globals.wgsl b/tests/out/wgsl/globals.wgsl index 7c9ca8056..73b436ddd 100644 --- a/tests/out/wgsl/globals.wgsl +++ b/tests/out/wgsl/globals.wgsl @@ -13,6 +13,8 @@ var alignment: Foo; var dummy: array>; @group(0) @binding(3) var float_vecs: array,20>; +@group(0) @binding(4) +var global_vec: vec4; fn test_msl_packed_vec3_as_arg(arg: vec3) { return; @@ -24,8 +26,8 @@ fn test_msl_packed_vec3_() { alignment.v3_ = vec3(1.0); alignment.v3_.x = 1.0; alignment.v3_.x = 2.0; - let _e19 = idx; - alignment.v3_[_e19] = 3.0; + let _e20 = idx; + alignment.v3_[_e20] = 3.0; let data = alignment; let unnamed = data.v3_; let unnamed_1 = data.v3_.zx; @@ -42,10 +44,16 @@ fn main() { var at: bool = true; test_msl_packed_vec3_(); - let _e9 = alignment.v1_; - wg[3] = _e9; - let _e14 = alignment.v3_.x; - wg[2] = _e14; + let _e10 = global_vec.x; + wg[6] = _e10; + let _e16 = dummy[1].y; + wg[5] = _e16; + let _e22 = float_vecs[0].w; + wg[4] = _e22; + let _e26 = alignment.v1_; + wg[3] = _e26; + let _e31 = alignment.v3_.x; + wg[2] = _e31; alignment.v1_ = 4.0; wg[1] = f32(arrayLength((&dummy))); atomicStore((&at_1), 2u);