From b95a72b2f6c209e7a3a40adfd90ffa281f47360f Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Fri, 29 Sep 2023 20:09:59 +0200 Subject: [PATCH] [wgsl] test `@workgroup_size` attribute with constants --- tests/in/const-exprs.wgsl | 5 +- tests/out/glsl/const-exprs.main.Compute.glsl | 4 +- tests/out/hlsl/const-exprs.hlsl | 4 +- tests/out/msl/const-exprs.msl | 2 + tests/out/spv/const-exprs.spvasm | 195 ++++++++++--------- tests/out/wgsl/const-exprs.wgsl | 4 +- 6 files changed, 113 insertions(+), 101 deletions(-) diff --git a/tests/in/const-exprs.wgsl b/tests/in/const-exprs.wgsl index 49e95a275f..11fee4a41c 100644 --- a/tests/in/const-exprs.wgsl +++ b/tests/in/const-exprs.wgsl @@ -1,7 +1,10 @@ @group(0) @binding(0) var out: vec4; @group(0) @binding(1) var out2: i32; -@compute @workgroup_size(1) +const TWO: u32 = 2u; +const THREE: i32 = 3i; + +@compute @workgroup_size(TWO, THREE, TWO - 1u) fn main() { swizzle_of_compose(); index_of_compose(); diff --git a/tests/out/glsl/const-exprs.main.Compute.glsl b/tests/out/glsl/const-exprs.main.Compute.glsl index e75ffa6506..8a1c3d45c7 100644 --- a/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/tests/out/glsl/const-exprs.main.Compute.glsl @@ -3,8 +3,10 @@ precision highp float; precision highp int; -layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +layout(local_size_x = 2, local_size_y = 3, local_size_z = 1) in; +const uint TWO = 2u; +const int THREE = 3; const int FOUR = 4; const int FOUR_ALIAS = 4; const int TEST_CONSTANT_ADDITION = 8; diff --git a/tests/out/hlsl/const-exprs.hlsl b/tests/out/hlsl/const-exprs.hlsl index 91b82ec56d..08cf89091e 100644 --- a/tests/out/hlsl/const-exprs.hlsl +++ b/tests/out/hlsl/const-exprs.hlsl @@ -1,3 +1,5 @@ +static const uint TWO = 2u; +static const int THREE = 3; static const int FOUR = 4; static const int FOUR_ALIAS = 4; static const int TEST_CONSTANT_ADDITION = 8; @@ -58,7 +60,7 @@ void compose_of_constant() return; } -[numthreads(1, 1, 1)] +[numthreads(2, 3, 1)] void main() { swizzle_of_compose(); diff --git a/tests/out/msl/const-exprs.msl b/tests/out/msl/const-exprs.msl index 5cbc23a006..4e6d67f801 100644 --- a/tests/out/msl/const-exprs.msl +++ b/tests/out/msl/const-exprs.msl @@ -4,6 +4,8 @@ using metal::uint; +constant uint TWO = 2u; +constant int THREE = 3; constant int FOUR = 4; constant int FOUR_ALIAS = 4; constant int TEST_CONSTANT_ADDITION = 8; diff --git a/tests/out/spv/const-exprs.spvasm b/tests/out/spv/const-exprs.spvasm index 1f7bf5c8ce..c8d41c59fa 100644 --- a/tests/out/spv/const-exprs.spvasm +++ b/tests/out/spv/const-exprs.spvasm @@ -1,127 +1,128 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 82 +; Bound: 83 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %72 "main" -OpExecutionMode %72 LocalSize 1 1 1 -OpDecorate %7 DescriptorSet 0 -OpDecorate %7 Binding 0 -OpDecorate %8 Block -OpMemberDecorate %8 0 Offset 0 +OpEntryPoint GLCompute %73 "main" +OpExecutionMode %73 LocalSize 2 3 1 OpDecorate %10 DescriptorSet 0 -OpDecorate %10 Binding 1 +OpDecorate %10 Binding 0 OpDecorate %11 Block OpMemberDecorate %11 0 Offset 0 +OpDecorate %13 DescriptorSet 0 +OpDecorate %13 Binding 1 +OpDecorate %14 Block +OpMemberDecorate %14 0 Offset 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpTypeVector %4 4 -%5 = OpConstant %4 4 -%6 = OpConstant %4 8 -%8 = OpTypeStruct %3 -%9 = OpTypePointer StorageBuffer %8 -%7 = OpVariable %9 StorageBuffer -%11 = OpTypeStruct %4 +%5 = OpTypeInt 32 0 +%6 = OpConstant %5 2 +%7 = OpConstant %4 3 +%8 = OpConstant %4 4 +%9 = OpConstant %4 8 +%11 = OpTypeStruct %3 %12 = OpTypePointer StorageBuffer %11 %10 = OpVariable %12 StorageBuffer -%15 = OpTypeFunction %2 -%16 = OpTypePointer StorageBuffer %3 -%18 = OpTypeInt 32 0 -%17 = OpConstant %18 0 -%20 = OpConstant %4 1 -%21 = OpConstant %4 2 -%22 = OpConstant %4 3 -%23 = OpConstantComposite %3 %5 %22 %21 %20 -%27 = OpTypePointer StorageBuffer %4 -%35 = OpConstant %4 6 -%42 = OpConstant %4 30 -%43 = OpConstant %4 70 -%45 = OpTypePointer Function %4 -%47 = OpConstantNull %4 -%49 = OpConstantNull %4 -%64 = OpConstant %4 -4 -%65 = OpConstantComposite %3 %64 %64 %64 %64 -%14 = OpFunction %2 None %15 -%13 = OpLabel -%19 = OpAccessChain %16 %7 %17 -OpBranch %24 -%24 = OpLabel -OpStore %19 %23 -OpReturn -OpFunctionEnd -%26 = OpFunction %2 None %15 +%14 = OpTypeStruct %4 +%15 = OpTypePointer StorageBuffer %14 +%13 = OpVariable %15 StorageBuffer +%18 = OpTypeFunction %2 +%19 = OpTypePointer StorageBuffer %3 +%20 = OpConstant %5 0 +%22 = OpConstant %4 1 +%23 = OpConstant %4 2 +%24 = OpConstantComposite %3 %8 %7 %23 %22 +%28 = OpTypePointer StorageBuffer %4 +%36 = OpConstant %4 6 +%43 = OpConstant %4 30 +%44 = OpConstant %4 70 +%46 = OpTypePointer Function %4 +%48 = OpConstantNull %4 +%50 = OpConstantNull %4 +%65 = OpConstant %4 -4 +%66 = OpConstantComposite %3 %65 %65 %65 %65 +%17 = OpFunction %2 None %18 +%16 = OpLabel +%21 = OpAccessChain %19 %10 %20 +OpBranch %25 %25 = OpLabel -%28 = OpAccessChain %27 %10 %17 -OpBranch %29 -%29 = OpLabel -%30 = OpLoad %4 %28 -%31 = OpIAdd %4 %30 %21 -OpStore %28 %31 +OpStore %21 %24 OpReturn OpFunctionEnd -%33 = OpFunction %2 None %15 -%32 = OpLabel -%34 = OpAccessChain %27 %10 %17 -OpBranch %36 -%36 = OpLabel -%37 = OpLoad %4 %34 -%38 = OpIAdd %4 %37 %35 -OpStore %34 %38 +%27 = OpFunction %2 None %18 +%26 = OpLabel +%29 = OpAccessChain %28 %13 %20 +OpBranch %30 +%30 = OpLabel +%31 = OpLoad %4 %29 +%32 = OpIAdd %4 %31 %23 +OpStore %29 %32 OpReturn OpFunctionEnd -%40 = OpFunction %2 None %15 -%39 = OpLabel -%46 = OpVariable %45 Function %47 -%50 = OpVariable %45 Function %43 -%44 = OpVariable %45 Function %42 -%48 = OpVariable %45 Function %49 -%41 = OpAccessChain %16 %7 %17 -OpBranch %51 -%51 = OpLabel -%52 = OpLoad %4 %44 -OpStore %46 %52 -%53 = OpLoad %4 %46 -OpStore %48 %53 -%54 = OpLoad %4 %44 -%55 = OpLoad %4 %46 -%56 = OpLoad %4 %48 -%57 = OpLoad %4 %50 -%58 = OpCompositeConstruct %3 %54 %55 %56 %57 -%59 = OpLoad %3 %41 -%60 = OpIAdd %3 %59 %58 -OpStore %41 %60 +%34 = OpFunction %2 None %18 +%33 = OpLabel +%35 = OpAccessChain %28 %13 %20 +OpBranch %37 +%37 = OpLabel +%38 = OpLoad %4 %35 +%39 = OpIAdd %4 %38 %36 +OpStore %35 %39 OpReturn OpFunctionEnd -%62 = OpFunction %2 None %15 -%61 = OpLabel -%63 = OpAccessChain %16 %7 %17 -OpBranch %66 -%66 = OpLabel -OpStore %63 %65 +%41 = OpFunction %2 None %18 +%40 = OpLabel +%47 = OpVariable %46 Function %48 +%51 = OpVariable %46 Function %44 +%45 = OpVariable %46 Function %43 +%49 = OpVariable %46 Function %50 +%42 = OpAccessChain %19 %10 %20 +OpBranch %52 +%52 = OpLabel +%53 = OpLoad %4 %45 +OpStore %47 %53 +%54 = OpLoad %4 %47 +OpStore %49 %54 +%55 = OpLoad %4 %45 +%56 = OpLoad %4 %47 +%57 = OpLoad %4 %49 +%58 = OpLoad %4 %51 +%59 = OpCompositeConstruct %3 %55 %56 %57 %58 +%60 = OpLoad %3 %42 +%61 = OpIAdd %3 %60 %59 +OpStore %42 %61 OpReturn OpFunctionEnd -%68 = OpFunction %2 None %15 +%63 = OpFunction %2 None %18 +%62 = OpLabel +%64 = OpAccessChain %19 %10 %20 +OpBranch %67 %67 = OpLabel -%69 = OpAccessChain %16 %7 %17 -OpBranch %70 -%70 = OpLabel -OpStore %69 %65 +OpStore %64 %66 OpReturn OpFunctionEnd -%72 = OpFunction %2 None %15 +%69 = OpFunction %2 None %18 +%68 = OpLabel +%70 = OpAccessChain %19 %10 %20 +OpBranch %71 %71 = OpLabel -%73 = OpAccessChain %16 %7 %17 -%74 = OpAccessChain %27 %10 %17 -OpBranch %75 -%75 = OpLabel -%76 = OpFunctionCall %2 %14 -%77 = OpFunctionCall %2 %26 -%78 = OpFunctionCall %2 %33 -%79 = OpFunctionCall %2 %40 -%80 = OpFunctionCall %2 %62 -%81 = OpFunctionCall %2 %68 +OpStore %70 %66 +OpReturn +OpFunctionEnd +%73 = OpFunction %2 None %18 +%72 = OpLabel +%74 = OpAccessChain %19 %10 %20 +%75 = OpAccessChain %28 %13 %20 +OpBranch %76 +%76 = OpLabel +%77 = OpFunctionCall %2 %17 +%78 = OpFunctionCall %2 %27 +%79 = OpFunctionCall %2 %34 +%80 = OpFunctionCall %2 %41 +%81 = OpFunctionCall %2 %63 +%82 = OpFunctionCall %2 %69 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/const-exprs.wgsl b/tests/out/wgsl/const-exprs.wgsl index 2ddc903aad..df675a8962 100644 --- a/tests/out/wgsl/const-exprs.wgsl +++ b/tests/out/wgsl/const-exprs.wgsl @@ -1,3 +1,5 @@ +const TWO: u32 = 2u; +const THREE: i32 = 3; const FOUR: i32 = 4; const FOUR_ALIAS: i32 = 4; const TEST_CONSTANT_ADDITION: i32 = 8; @@ -54,7 +56,7 @@ fn compose_of_constant() { return; } -@compute @workgroup_size(1, 1, 1) +@compute @workgroup_size(2, 3, 1) fn main() { swizzle_of_compose(); index_of_compose();