From bdbf821c3247742f0fa069636d03937b8166c7aa Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 12 Apr 2022 19:02:37 -0700 Subject: [PATCH] Test matrix-typed uniform globals. --- tests/in/globals.wgsl | 5 +- tests/out/glsl/globals.main.Compute.glsl | 27 ++- tests/out/hlsl/globals.hlsl | 26 ++- tests/out/msl/globals.msl | 32 +-- tests/out/spv/globals.spvasm | 278 ++++++++++++----------- tests/out/wgsl/globals.wgsl | 27 ++- 6 files changed, 211 insertions(+), 184 deletions(-) diff --git a/tests/in/globals.wgsl b/tests/in/globals.wgsl index 573c1575c7..a4467eae75 100644 --- a/tests/in/globals.wgsl +++ b/tests/in/globals.wgsl @@ -22,6 +22,9 @@ var float_vecs: array, 20>; @group(0) @binding(4) var global_vec: vec4; +@group(0) @binding(5) +var global_mat: mat4x4; + fn test_msl_packed_vec3_as_arg(arg: vec3) {} fn test_msl_packed_vec3() { @@ -53,7 +56,7 @@ fn test_msl_packed_vec3() { fn main() { test_msl_packed_vec3(); - wg[6] = global_vec.x; + wg[6] = (global_mat * global_vec).x; wg[5] = dummy[1].y; wg[4] = float_vecs[0].w; wg[3] = alignment.v1; diff --git a/tests/out/glsl/globals.main.Compute.glsl b/tests/out/glsl/globals.main.Compute.glsl index 785a3d1e1e..40066567f3 100644 --- a/tests/out/glsl/globals.main.Compute.glsl +++ b/tests/out/glsl/globals.main.Compute.glsl @@ -21,6 +21,8 @@ uniform type_8_block_2Compute { vec4 _group_0_binding_3_cs[20]; }; uniform type_7_block_3Compute { vec4 _group_0_binding_4_cs; }; +uniform type_9_block_4Compute { mat4x4 _group_0_binding_5_cs; }; + void test_msl_packed_vec3_as_arg(vec3 arg) { return; @@ -31,8 +33,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 _e20 = idx; - _group_0_binding_1_cs.v3_[_e20] = 3.0; + int _e21 = idx; + _group_0_binding_1_cs.v3_[_e21] = 3.0; Foo data = _group_0_binding_1_cs; vec3 unnamed = data.v3_; vec2 unnamed_1 = data.v3_.zx; @@ -47,16 +49,17 @@ void main() { float Foo_1 = 1.0; bool at = true; test_msl_packed_vec3_(); - 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; + mat4x4 _e10 = _group_0_binding_5_cs; + vec4 _e11 = _group_0_binding_4_cs; + wg[6] = (_e10 * _e11).x; + float _e19 = _group_0_binding_2_cs[1].y; + wg[5] = _e19; + float _e25 = _group_0_binding_3_cs[0].w; + wg[4] = _e25; + float _e29 = _group_0_binding_1_cs.v1_; + wg[3] = _e29; + float _e34 = _group_0_binding_1_cs.v3_.x; + wg[2] = _e34; _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 abf543ff65..cb9e1b4bb8 100644 --- a/tests/out/hlsl/globals.hlsl +++ b/tests/out/hlsl/globals.hlsl @@ -11,6 +11,7 @@ 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; } +cbuffer global_mat : register(b5) { float4x4 global_mat; } void test_msl_packed_vec3_as_arg(float3 arg) { @@ -24,8 +25,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 _expr20 = idx; - alignment.Store(_expr20*4+0, asuint(3.0)); + int _expr21 = idx; + alignment.Store(_expr21*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; @@ -50,16 +51,17 @@ void main() bool at = true; test_msl_packed_vec3_(); - 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; + float4x4 _expr10 = global_mat; + float4 _expr11 = global_vec; + wg[6] = mul(_expr11, _expr10).x; + float _expr19 = asfloat(dummy.Load(4+8)); + wg[5] = _expr19; + float _expr25 = float_vecs[0].w; + wg[4] = _expr25; + float _expr29 = asfloat(alignment.Load(12)); + wg[3] = _expr29; + float _expr34 = asfloat(alignment.Load(0+0)); + wg[2] = _expr34; 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 3c83f67dd2..b02d6056d4 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -21,7 +21,7 @@ struct type_8 { metal::float4 inner[20]; }; constant metal::float3 const_type_4_ = {0.0, 0.0, 0.0}; -constant metal::float3x3 const_type_10_ = {const_type_4_, const_type_4_, const_type_4_}; +constant metal::float3x3 const_type_11_ = {const_type_4_, const_type_4_, const_type_4_}; void test_msl_packed_vec3_as_arg( metal::float3 arg @@ -36,14 +36,14 @@ void test_msl_packed_vec3_( alignment.v3_ = metal::float3(1.0); alignment.v3_[0] = 1.0; alignment.v3_[0] = 2.0; - int _e20 = idx; - alignment.v3_[_e20] = 3.0; + int _e21 = idx; + alignment.v3_[_e21] = 3.0; Foo data = alignment; metal::float3 unnamed = data.v3_; metal::float2 unnamed_1 = metal::float3(data.v3_).zx; test_msl_packed_vec3_as_arg(data.v3_); - metal::float3 unnamed_2 = metal::float3(data.v3_) * const_type_10_; - metal::float3 unnamed_3 = const_type_10_ * metal::float3(data.v3_); + metal::float3 unnamed_2 = metal::float3(data.v3_) * const_type_11_; + metal::float3 unnamed_3 = const_type_11_ * metal::float3(data.v3_); metal::float3 unnamed_4 = data.v3_ * 2.0; metal::float3 unnamed_5 = 2.0 * data.v3_; } @@ -55,21 +55,23 @@ kernel void main_( , device type_6 const& dummy [[user(fake0)]] , constant type_8& float_vecs [[user(fake0)]] , constant metal::float4& global_vec [[user(fake0)]] +, constant metal::float4x4& global_mat [[user(fake0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { float Foo_1 = 1.0; bool at = true; test_msl_packed_vec3_(alignment); - 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; + metal::float4x4 _e10 = global_mat; + metal::float4 _e11 = global_vec; + wg.inner[6] = (_e10 * _e11).x; + float _e19 = dummy[1].y; + wg.inner[5] = _e19; + float _e25 = float_vecs.inner[0].w; + wg.inner[4] = _e25; + float _e29 = alignment.v1_; + wg.inner[3] = _e29; + float _e34 = alignment.v3_[0]; + wg.inner[2] = _e34; 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 07e6da2a77..8109c3ebfd 100644 --- a/tests/out/spv/globals.spvasm +++ b/tests/out/spv/globals.spvasm @@ -1,35 +1,41 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 130 +; Bound: 138 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %94 "main" -OpExecutionMode %94 LocalSize 1 1 1 +OpEntryPoint GLCompute %99 "main" +OpExecutionMode %99 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 1 -OpDecorate %39 Block -OpMemberDecorate %39 0 Offset 0 -OpDecorate %41 NonWritable -OpDecorate %41 DescriptorSet 0 -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 +OpDecorate %39 DescriptorSet 0 +OpDecorate %39 Binding 1 +OpDecorate %40 Block +OpMemberDecorate %40 0 Offset 0 +OpDecorate %42 NonWritable +OpDecorate %42 DescriptorSet 0 +OpDecorate %42 Binding 2 +OpDecorate %43 Block +OpMemberDecorate %43 0 Offset 0 +OpDecorate %45 DescriptorSet 0 +OpDecorate %45 Binding 3 +OpDecorate %46 Block +OpMemberDecorate %46 0 Offset 0 +OpDecorate %48 DescriptorSet 0 +OpDecorate %48 Binding 4 +OpDecorate %49 Block +OpMemberDecorate %49 0 Offset 0 +OpDecorate %51 DescriptorSet 0 +OpDecorate %51 Binding 5 +OpDecorate %52 Block +OpMemberDecorate %52 0 Offset 0 +OpMemberDecorate %52 0 ColMajor +OpMemberDecorate %52 0 MatrixStride 16 %2 = OpTypeVoid %4 = OpTypeBool %3 = OpConstantTrue %4 @@ -59,123 +65,131 @@ OpMemberDecorate %48 0 Offset 0 %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 %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 +%31 = OpTypeMatrix %29 4 +%32 = OpTypeMatrix %25 3 +%33 = OpConstantComposite %25 %15 %15 %15 +%34 = OpConstantComposite %32 %33 %33 %33 +%36 = OpTypePointer Workgroup %24 +%35 = OpVariable %36 Workgroup +%38 = OpTypePointer Workgroup %6 +%37 = OpVariable %38 Workgroup +%40 = OpTypeStruct %26 +%41 = OpTypePointer StorageBuffer %40 +%39 = OpVariable %41 StorageBuffer +%43 = OpTypeStruct %28 +%44 = OpTypePointer StorageBuffer %43 +%42 = OpVariable %44 StorageBuffer +%46 = OpTypeStruct %30 +%47 = OpTypePointer Uniform %46 +%45 = OpVariable %47 Uniform +%49 = OpTypeStruct %29 +%50 = OpTypePointer Uniform %49 +%48 = OpVariable %50 Uniform +%52 = OpTypeStruct %31 +%53 = OpTypePointer Uniform %52 +%51 = OpVariable %53 Uniform +%57 = OpTypeFunction %2 %25 +%58 = OpTypePointer StorageBuffer %28 +%59 = OpTypePointer Uniform %29 +%60 = OpTypePointer StorageBuffer %26 +%61 = OpTypePointer Uniform %30 +%62 = OpTypePointer Uniform %31 +%65 = OpTypePointer Function %8 +%68 = OpTypeFunction %2 +%69 = OpConstant %6 0 +%72 = OpTypePointer StorageBuffer %25 +%75 = OpTypePointer StorageBuffer %10 +%95 = OpTypePointer Function %10 +%97 = OpTypePointer Function %4 +%107 = OpTypePointer Workgroup %10 +%112 = OpConstant %6 6 +%114 = OpTypePointer StorageBuffer %27 +%115 = OpConstant %6 1 +%118 = OpConstant %6 5 +%120 = OpTypePointer Uniform %10 +%121 = OpConstant %6 3 +%124 = OpConstant %6 4 +%126 = OpTypePointer StorageBuffer %10 +%137 = OpConstant %6 256 +%56 = OpFunction %2 None %57 +%55 = OpFunctionParameter %25 +%54 = OpLabel +OpBranch %63 +%63 = OpLabel OpReturn OpFunctionEnd -%62 = OpFunction %2 None %63 -%61 = OpLabel -%59 = OpVariable %60 Function %11 -%65 = OpAccessChain %54 %38 %64 -OpBranch %66 +%67 = OpFunction %2 None %68 %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 +%64 = OpVariable %65 Function %11 +%70 = OpAccessChain %60 %39 %69 +OpBranch %71 +%71 = OpLabel +%73 = OpCompositeConstruct %25 %9 %9 %9 +%74 = OpAccessChain %72 %70 %69 +OpStore %74 %73 +%76 = OpAccessChain %75 %70 %69 %69 +OpStore %76 %9 +%77 = OpAccessChain %75 %70 %69 %69 +OpStore %77 %13 +%78 = OpLoad %8 %64 +%79 = OpAccessChain %75 %70 %69 %78 +OpStore %79 %14 +%80 = OpLoad %26 %70 +%81 = OpCompositeExtract %25 %80 0 +%82 = OpCompositeExtract %25 %80 0 +%83 = OpVectorShuffle %27 %82 %82 2 0 +%84 = OpCompositeExtract %25 %80 0 +%85 = OpFunctionCall %2 %56 %84 +%86 = OpCompositeExtract %25 %80 0 +%87 = OpVectorTimesMatrix %25 %86 %34 +%88 = OpCompositeExtract %25 %80 0 +%89 = OpMatrixTimesVector %25 %34 %88 +%90 = OpCompositeExtract %25 %80 0 +%91 = OpVectorTimesScalar %25 %90 %13 +%92 = OpCompositeExtract %25 %80 0 +%93 = OpVectorTimesScalar %25 %92 %13 OpReturn OpFunctionEnd -%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 +%99 = OpFunction %2 None %68 +%98 = OpLabel +%94 = OpVariable %95 Function %9 +%96 = OpVariable %97 Function %23 +%100 = OpAccessChain %60 %39 %69 +%101 = OpAccessChain %58 %42 %69 +%102 = OpAccessChain %61 %45 %69 +%103 = OpAccessChain %59 %48 %69 +%104 = OpAccessChain %62 %51 %69 +OpBranch %105 +%105 = OpLabel +%106 = OpFunctionCall %2 %67 +%108 = OpLoad %31 %104 +%109 = OpLoad %29 %103 +%110 = OpMatrixTimesVector %29 %108 %109 +%111 = OpCompositeExtract %10 %110 0 +%113 = OpAccessChain %107 %35 %112 +OpStore %113 %111 +%116 = OpAccessChain %75 %101 %115 %115 +%117 = OpLoad %10 %116 +%119 = OpAccessChain %107 %35 %118 +OpStore %119 %117 +%122 = OpAccessChain %120 %102 %69 %121 %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 +%125 = OpAccessChain %107 %35 %124 +OpStore %125 %123 +%127 = OpAccessChain %126 %100 %115 +%128 = OpLoad %10 %127 +%129 = OpAccessChain %107 %35 %121 +OpStore %129 %128 +%130 = OpAccessChain %75 %100 %69 %69 +%131 = OpLoad %10 %130 +%132 = OpAccessChain %107 %35 %22 +OpStore %132 %131 +%133 = OpAccessChain %126 %100 %115 +OpStore %133 %21 +%134 = OpArrayLength %6 %42 0 +%135 = OpConvertUToF %10 %134 +%136 = OpAccessChain %107 %35 %115 +OpStore %136 %135 +OpAtomicStore %37 %20 %137 %22 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/globals.wgsl b/tests/out/wgsl/globals.wgsl index 73b436ddd3..0060e40bf5 100644 --- a/tests/out/wgsl/globals.wgsl +++ b/tests/out/wgsl/globals.wgsl @@ -15,6 +15,8 @@ var dummy: array>; var float_vecs: array,20>; @group(0) @binding(4) var global_vec: vec4; +@group(0) @binding(5) +var global_mat: mat4x4; fn test_msl_packed_vec3_as_arg(arg: vec3) { return; @@ -26,8 +28,8 @@ fn test_msl_packed_vec3_() { alignment.v3_ = vec3(1.0); alignment.v3_.x = 1.0; alignment.v3_.x = 2.0; - let _e20 = idx; - alignment.v3_[_e20] = 3.0; + let _e21 = idx; + alignment.v3_[_e21] = 3.0; let data = alignment; let unnamed = data.v3_; let unnamed_1 = data.v3_.zx; @@ -44,16 +46,17 @@ fn main() { var at: bool = true; test_msl_packed_vec3_(); - 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; + let _e10 = global_mat; + let _e11 = global_vec; + wg[6] = (_e10 * _e11).x; + let _e19 = dummy[1].y; + wg[5] = _e19; + let _e25 = float_vecs[0].w; + wg[4] = _e25; + let _e29 = alignment.v1_; + wg[3] = _e29; + let _e34 = alignment.v3_.x; + wg[2] = _e34; alignment.v1_ = 4.0; wg[1] = f32(arrayLength((&dummy))); atomicStore((&at_1), 2u);