diff --git a/tests/in/const-exprs.wgsl b/tests/in/const-exprs.wgsl index fd1b9c31bf..73d5b4af2f 100644 --- a/tests/in/const-exprs.wgsl +++ b/tests/in/const-exprs.wgsl @@ -7,6 +7,8 @@ fn main() { index_of_compose(); compose_three_deep(); non_constant_initializers(); + splat_of_constant(); + compose_of_constant(); } // Swizzle the value of nested Compose expressions. @@ -46,3 +48,15 @@ fn non_constant_initializers() { out += vec4(w, x, y, z); } + +// Constant evaluation should be able to see through constants to +// their values. +const FOUR: i32 = 4; + +fn splat_of_constant() { + out = -vec4(FOUR); +} + +fn compose_of_constant() { + out = -vec4(FOUR, FOUR, FOUR, FOUR); +} diff --git a/tests/out/glsl/const-exprs.main.Compute.glsl b/tests/out/glsl/const-exprs.main.Compute.glsl index df029bdaaa..fde4c96156 100644 --- a/tests/out/glsl/const-exprs.main.Compute.glsl +++ b/tests/out/glsl/const-exprs.main.Compute.glsl @@ -5,6 +5,8 @@ precision highp int; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +const int FOUR = 4; + layout(std430) buffer type_block_0Compute { ivec4 _group_0_binding_0_cs; }; layout(std430) buffer type_1_block_1Compute { int _group_0_binding_1_cs; }; @@ -49,11 +51,23 @@ void non_constant_initializers() { return; } +void splat_of_constant() { + _group_0_binding_0_cs = -(ivec4(FOUR)); + return; +} + +void compose_of_constant() { + _group_0_binding_0_cs = -(ivec4(FOUR, FOUR, FOUR, FOUR)); + return; +} + void main() { swizzle_of_compose(); index_of_compose(); compose_three_deep(); non_constant_initializers(); + splat_of_constant(); + compose_of_constant(); return; } diff --git a/tests/out/hlsl/const-exprs.hlsl b/tests/out/hlsl/const-exprs.hlsl index a43f27b02a..6e325a9bf1 100644 --- a/tests/out/hlsl/const-exprs.hlsl +++ b/tests/out/hlsl/const-exprs.hlsl @@ -1,3 +1,5 @@ +static const int FOUR = 4; + RWByteAddressBuffer out_ : register(u0); RWByteAddressBuffer out2_ : register(u1); @@ -45,6 +47,18 @@ void non_constant_initializers() return; } +void splat_of_constant() +{ + out_.Store4(0, asuint(-((FOUR).xxxx))); + return; +} + +void compose_of_constant() +{ + out_.Store4(0, asuint(-(int4(FOUR, FOUR, FOUR, FOUR)))); + return; +} + [numthreads(1, 1, 1)] void main() { @@ -52,5 +66,7 @@ void main() index_of_compose(); compose_three_deep(); non_constant_initializers(); + splat_of_constant(); + compose_of_constant(); return; } diff --git a/tests/out/msl/const-exprs.msl b/tests/out/msl/const-exprs.msl index a9aa559993..1d9c8e9caf 100644 --- a/tests/out/msl/const-exprs.msl +++ b/tests/out/msl/const-exprs.msl @@ -4,6 +4,7 @@ using metal::uint; +constant int FOUR = 4; void swizzle_of_compose( device metal::int4& out @@ -52,6 +53,20 @@ void non_constant_initializers( return; } +void splat_of_constant( + device metal::int4& out +) { + out = -(metal::int4(FOUR)); + return; +} + +void compose_of_constant( + device metal::int4& out +) { + out = -(metal::int4(FOUR, FOUR, FOUR, FOUR)); + return; +} + kernel void main_( device metal::int4& out [[user(fake0)]] , device int& out2_ [[user(fake0)]] @@ -60,5 +75,7 @@ kernel void main_( index_of_compose(out2_); compose_three_deep(out2_); non_constant_initializers(out); + splat_of_constant(out); + compose_of_constant(out); return; } diff --git a/tests/out/spv/const-exprs.spvasm b/tests/out/spv/const-exprs.spvasm index 50b0373a46..9e0d6c14cf 100644 --- a/tests/out/spv/const-exprs.spvasm +++ b/tests/out/spv/const-exprs.spvasm @@ -1,42 +1,42 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 72 +; Bound: 86 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %64 "main" -OpExecutionMode %64 LocalSize 1 1 1 -OpDecorate %6 DescriptorSet 0 -OpDecorate %6 Binding 0 -OpDecorate %7 Block -OpMemberDecorate %7 0 Offset 0 -OpDecorate %9 DescriptorSet 0 -OpDecorate %9 Binding 1 -OpDecorate %10 Block -OpMemberDecorate %10 0 Offset 0 +OpEntryPoint GLCompute %76 "main" +OpExecutionMode %76 LocalSize 1 1 1 +OpDecorate %7 DescriptorSet 0 +OpDecorate %7 Binding 0 +OpDecorate %8 Block +OpMemberDecorate %8 0 Offset 0 +OpDecorate %10 DescriptorSet 0 +OpDecorate %10 Binding 1 +OpDecorate %11 Block +OpMemberDecorate %11 0 Offset 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpTypeVector %4 4 %5 = OpTypeVector %4 2 -%7 = OpTypeStruct %3 -%8 = OpTypePointer StorageBuffer %7 -%6 = OpVariable %8 StorageBuffer -%10 = OpTypeStruct %4 -%11 = OpTypePointer StorageBuffer %10 -%9 = OpVariable %11 StorageBuffer -%14 = OpTypeFunction %2 -%15 = OpTypePointer StorageBuffer %3 -%17 = OpTypeInt 32 0 -%16 = OpConstant %17 0 -%19 = OpConstant %4 1 -%20 = OpConstant %4 2 -%21 = OpConstantComposite %5 %19 %20 -%22 = OpConstant %4 3 -%23 = OpConstant %4 4 -%24 = OpConstantComposite %5 %22 %23 -%25 = OpConstantComposite %3 %23 %22 %20 %19 +%6 = OpConstant %4 4 +%8 = OpTypeStruct %3 +%9 = OpTypePointer StorageBuffer %8 +%7 = OpVariable %9 StorageBuffer +%11 = OpTypeStruct %4 +%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 = OpConstantComposite %5 %20 %21 +%23 = OpConstant %4 3 +%24 = OpConstantComposite %5 %23 %6 +%25 = OpConstantComposite %3 %6 %23 %21 %20 %29 = OpTypePointer StorageBuffer %4 %37 = OpConstant %4 6 %44 = OpConstant %4 30 @@ -44,27 +44,29 @@ OpMemberDecorate %10 0 Offset 0 %47 = OpTypePointer Function %4 %49 = OpConstantNull %4 %51 = OpConstantNull %4 -%13 = OpFunction %2 None %14 -%12 = OpLabel -%18 = OpAccessChain %15 %6 %16 +%66 = OpConstantComposite %3 %6 %6 %6 %6 +%72 = OpConstantComposite %3 %6 %6 %6 %6 +%14 = OpFunction %2 None %15 +%13 = OpLabel +%19 = OpAccessChain %16 %7 %17 OpBranch %26 %26 = OpLabel -OpStore %18 %25 +OpStore %19 %25 OpReturn OpFunctionEnd -%28 = OpFunction %2 None %14 +%28 = OpFunction %2 None %15 %27 = OpLabel -%30 = OpAccessChain %29 %9 %16 +%30 = OpAccessChain %29 %10 %17 OpBranch %31 %31 = OpLabel %32 = OpLoad %4 %30 -%33 = OpIAdd %4 %32 %20 +%33 = OpIAdd %4 %32 %21 OpStore %30 %33 OpReturn OpFunctionEnd -%35 = OpFunction %2 None %14 +%35 = OpFunction %2 None %15 %34 = OpLabel -%36 = OpAccessChain %29 %9 %16 +%36 = OpAccessChain %29 %10 %17 OpBranch %38 %38 = OpLabel %39 = OpLoad %4 %36 @@ -72,13 +74,13 @@ OpBranch %38 OpStore %36 %40 OpReturn OpFunctionEnd -%42 = OpFunction %2 None %14 +%42 = OpFunction %2 None %15 %41 = OpLabel %48 = OpVariable %47 Function %49 %52 = OpVariable %47 Function %45 %46 = OpVariable %47 Function %44 %50 = OpVariable %47 Function %51 -%43 = OpAccessChain %15 %6 %16 +%43 = OpAccessChain %16 %7 %17 OpBranch %53 %53 = OpLabel %54 = OpLoad %4 %46 @@ -95,15 +97,35 @@ OpStore %50 %55 OpStore %43 %62 OpReturn OpFunctionEnd -%64 = OpFunction %2 None %14 +%64 = OpFunction %2 None %15 %63 = OpLabel -%65 = OpAccessChain %15 %6 %16 -%66 = OpAccessChain %29 %9 %16 +%65 = OpAccessChain %16 %7 %17 OpBranch %67 %67 = OpLabel -%68 = OpFunctionCall %2 %13 -%69 = OpFunctionCall %2 %28 -%70 = OpFunctionCall %2 %35 -%71 = OpFunctionCall %2 %42 +%68 = OpSNegate %3 %66 +OpStore %65 %68 +OpReturn +OpFunctionEnd +%70 = OpFunction %2 None %15 +%69 = OpLabel +%71 = OpAccessChain %16 %7 %17 +OpBranch %73 +%73 = OpLabel +%74 = OpSNegate %3 %72 +OpStore %71 %74 +OpReturn +OpFunctionEnd +%76 = OpFunction %2 None %15 +%75 = OpLabel +%77 = OpAccessChain %16 %7 %17 +%78 = OpAccessChain %29 %10 %17 +OpBranch %79 +%79 = OpLabel +%80 = OpFunctionCall %2 %14 +%81 = OpFunctionCall %2 %28 +%82 = OpFunctionCall %2 %35 +%83 = OpFunctionCall %2 %42 +%84 = OpFunctionCall %2 %64 +%85 = OpFunctionCall %2 %70 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 dcef1812b9..47744ede0d 100644 --- a/tests/out/wgsl/const-exprs.wgsl +++ b/tests/out/wgsl/const-exprs.wgsl @@ -1,3 +1,5 @@ +const FOUR: i32 = 4; + @group(0) @binding(0) var out: vec4; @group(0) @binding(1) @@ -43,11 +45,23 @@ fn non_constant_initializers() { return; } +fn splat_of_constant() { + out = -(vec4(FOUR)); + return; +} + +fn compose_of_constant() { + out = -(vec4(FOUR, FOUR, FOUR, FOUR)); + return; +} + @compute @workgroup_size(1, 1, 1) fn main() { swizzle_of_compose(); index_of_compose(); compose_three_deep(); non_constant_initializers(); + splat_of_constant(); + compose_of_constant(); return; }