[naga] Change operators.wgsl to use values that matter.

Change `splat` to take run-time values as arguments, to prevent
constant evaluation from boiling everything away.
This commit is contained in:
Jim Blandy 2023-12-11 12:51:40 -08:00 committed by Teodor Tanasoaia
parent 439bf3c1b4
commit 0e41cfe4ca
6 changed files with 436 additions and 417 deletions

View File

@ -21,9 +21,9 @@ fn builtins() -> vec4<f32> {
return vec4<f32>(vec4<i32>(s1) + v_i32_zero) + s2 + m1 + m2 + b1 + b2; return vec4<f32>(vec4<i32>(s1) + v_i32_zero) + s2 + m1 + m2 + b1 + b2;
} }
fn splat() -> vec4<f32> { fn splat(m: f32, n: i32) -> vec4<f32> {
let a = (1.0 + vec2<f32>(2.0) - 3.0) / 4.0; let a = (2.0 + vec2<f32>(m) - 4.0) / 8.0;
let b = vec4<i32>(5) % 2; let b = vec4<i32>(n) % 2;
return a.xyxy + vec4<f32>(b); return a.xyxy + vec4<f32>(b);
} }
@ -280,9 +280,9 @@ fn assignment() {
} }
@compute @workgroup_size(1) @compute @workgroup_size(1)
fn main() { fn main(@builtin(workgroup_id) id: vec3<u32>) {
builtins(); builtins();
splat(); splat(f32(id.x), i32(id.y));
bool_cast(v_f32_one.xyz); bool_cast(v_f32_one.xyz);
logical(); logical();

View File

@ -23,9 +23,9 @@ vec4 builtins() {
return (((((vec4((ivec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_); return (((((vec4((ivec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_);
} }
vec4 splat() { vec4 splat(float m, int n) {
vec2 a_2 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0)); vec2 a_2 = (((vec2(2.0) + vec2(m)) - vec2(4.0)) / vec2(8.0));
ivec4 b = (ivec4(5) % ivec4(2)); ivec4 b = (ivec4(n) % ivec4(2));
return (a_2.xyxy + vec4(b)); return (a_2.xyxy + vec4(b));
} }
@ -247,9 +247,10 @@ void negation_avoids_prefix_decrement() {
} }
void main() { void main() {
vec4 _e0 = builtins(); uvec3 id = gl_WorkGroupID;
vec4 _e1 = splat(); vec4 _e1 = builtins();
vec3 _e6 = bool_cast(vec3(1.0, 1.0, 1.0)); vec4 _e6 = splat(float(id.x), int(id.y));
vec3 _e11 = bool_cast(vec3(1.0, 1.0, 1.0));
logical(); logical();
arithmetic(); arithmetic();
bit(); bit();

View File

@ -16,10 +16,10 @@ float4 builtins()
return (((((float4(((s1_).xxxx + v_i32_zero)) + s2_) + m1_) + m2_) + (b1_).xxxx) + b2_); return (((((float4(((s1_).xxxx + v_i32_zero)) + s2_) + m1_) + m2_) + (b1_).xxxx) + b2_);
} }
float4 splat() float4 splat(float m, int n)
{ {
float2 a_2 = ((((1.0).xx + (2.0).xx) - (3.0).xx) / (4.0).xx); float2 a_2 = ((((2.0).xx + (m).xx) - (4.0).xx) / (8.0).xx);
int4 b = ((5).xxxx % (2).xxxx); int4 b = ((n).xxxx % (2).xxxx);
return (a_2.xyxy + float4(b)); return (a_2.xyxy + float4(b));
} }
@ -251,11 +251,11 @@ void negation_avoids_prefix_decrement()
} }
[numthreads(1, 1, 1)] [numthreads(1, 1, 1)]
void main() void main(uint3 id : SV_GroupID)
{ {
const float4 _e0 = builtins(); const float4 _e1 = builtins();
const float4 _e1 = splat(); const float4 _e6 = splat(float(id.x), int(id.y));
const float3 _e6 = bool_cast(float3(1.0, 1.0, 1.0)); const float3 _e11 = bool_cast(float3(1.0, 1.0, 1.0));
logical(); logical();
arithmetic(); arithmetic();
bit(); bit();

View File

@ -23,9 +23,11 @@ metal::float4 builtins(
} }
metal::float4 splat( metal::float4 splat(
float m,
int n
) { ) {
metal::float2 a_2 = ((metal::float2(1.0) + metal::float2(2.0)) - metal::float2(3.0)) / metal::float2(4.0); metal::float2 a_2 = ((metal::float2(2.0) + metal::float2(m)) - metal::float2(4.0)) / metal::float2(8.0);
metal::int4 b = metal::int4(5) % metal::int4(2); metal::int4 b = metal::int4(n) % metal::int4(2);
return a_2.xyxy + static_cast<metal::float4>(b); return a_2.xyxy + static_cast<metal::float4>(b);
} }
@ -255,11 +257,14 @@ void negation_avoids_prefix_decrement(
int p7_ = -(-(-(-(-(1))))); int p7_ = -(-(-(-(-(1)))));
} }
struct main_Input {
};
kernel void main_( kernel void main_(
metal::uint3 id [[threadgroup_position_in_grid]]
) { ) {
metal::float4 _e0 = builtins(); metal::float4 _e1 = builtins();
metal::float4 _e1 = splat(); metal::float4 _e6 = splat(static_cast<float>(id.x), static_cast<int>(id.y));
metal::float3 _e6 = bool_cast(metal::float3(1.0, 1.0, 1.0)); metal::float3 _e11 = bool_cast(metal::float3(1.0, 1.0, 1.0));
logical(); logical();
arithmetic(); arithmetic();
bit(); bit();

View File

@ -1,12 +1,13 @@
; SPIR-V ; SPIR-V
; Version: 1.1 ; Version: 1.1
; Generator: rspirv ; Generator: rspirv
; Bound: 377 ; Bound: 389
OpCapability Shader OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450" %1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %366 "main" OpEntryPoint GLCompute %374 "main" %371
OpExecutionMode %366 LocalSize 1 1 1 OpExecutionMode %374 LocalSize 1 1 1
OpDecorate %371 BuiltIn WorkgroupId
%2 = OpTypeVoid %2 = OpTypeVoid
%4 = OpTypeFloat 32 %4 = OpTypeFloat 32
%3 = OpTypeVector %4 4 %3 = OpTypeVector %4 4
@ -20,418 +21,430 @@ OpExecutionMode %366 LocalSize 1 1 1
%12 = OpTypeMatrix %10 4 %12 = OpTypeMatrix %10 4
%13 = OpTypeMatrix %3 3 %13 = OpTypeMatrix %3 3
%14 = OpTypeVector %6 3 %14 = OpTypeVector %6 3
%15 = OpConstant %4 1.0 %16 = OpTypeInt 32 0
%16 = OpConstantComposite %3 %15 %15 %15 %15 %15 = OpTypeVector %16 3
%17 = OpConstant %4 0.0 %17 = OpConstant %4 1.0
%18 = OpConstantComposite %3 %17 %17 %17 %17 %18 = OpConstantComposite %3 %17 %17 %17 %17
%19 = OpConstant %4 0.5 %19 = OpConstant %4 0.0
%20 = OpConstantComposite %3 %19 %19 %19 %19 %20 = OpConstantComposite %3 %19 %19 %19 %19
%21 = OpConstant %6 1 %21 = OpConstant %4 0.5
%22 = OpConstantComposite %5 %21 %21 %21 %21 %22 = OpConstantComposite %3 %21 %21 %21 %21
%25 = OpTypeFunction %3 %23 = OpConstant %6 1
%26 = OpConstantTrue %8 %24 = OpConstantComposite %5 %23 %23 %23 %23
%27 = OpConstant %6 0 %27 = OpTypeFunction %3
%28 = OpConstantFalse %8 %28 = OpConstantTrue %8
%29 = OpConstantComposite %7 %28 %28 %28 %28 %29 = OpConstant %6 0
%30 = OpConstant %4 0.1 %30 = OpConstantFalse %8
%31 = OpConstantComposite %5 %27 %27 %27 %27 %31 = OpConstantComposite %7 %30 %30 %30 %30
%53 = OpConstant %4 2.0 %32 = OpConstant %4 0.1
%54 = OpConstantComposite %9 %53 %53 %33 = OpConstantComposite %5 %29 %29 %29 %29
%55 = OpConstantComposite %9 %15 %15 %57 = OpTypeFunction %3 %4 %6
%56 = OpConstant %4 3.0 %58 = OpConstant %4 2.0
%57 = OpConstantComposite %9 %56 %56
%58 = OpConstant %4 4.0
%59 = OpConstantComposite %9 %58 %58 %59 = OpConstantComposite %9 %58 %58
%60 = OpConstant %6 5 %60 = OpConstant %4 4.0
%61 = OpConstantComposite %5 %60 %60 %60 %60 %61 = OpConstantComposite %9 %60 %60
%62 = OpConstant %6 2 %62 = OpConstant %4 8.0
%63 = OpConstantComposite %5 %62 %62 %62 %62 %63 = OpConstantComposite %9 %62 %62
%74 = OpTypeFunction %9 %64 = OpConstant %6 2
%76 = OpTypePointer Function %9 %65 = OpConstantComposite %5 %64 %64 %64 %64
%88 = OpTypeFunction %10 %10 %78 = OpTypeFunction %9
%90 = OpTypeVector %8 3 %79 = OpConstantComposite %9 %17 %17
%91 = OpConstantComposite %10 %17 %17 %17 %80 = OpConstant %4 3.0
%93 = OpConstantComposite %10 %15 %15 %15 %81 = OpConstantComposite %9 %80 %80
%97 = OpTypeFunction %2 %83 = OpTypePointer Function %9
%98 = OpTypeVector %8 2 %95 = OpTypeFunction %10 %10
%99 = OpConstantComposite %98 %26 %26 %97 = OpTypeVector %8 3
%100 = OpConstantComposite %90 %26 %26 %26 %98 = OpConstantComposite %10 %19 %19 %19
%101 = OpConstantComposite %90 %28 %28 %28 %100 = OpConstantComposite %10 %17 %17 %17
%102 = OpConstantComposite %7 %26 %26 %26 %26 %104 = OpTypeFunction %2
%103 = OpConstantComposite %7 %28 %28 %28 %28 %105 = OpTypeVector %8 2
%115 = OpTypeInt 32 0 %106 = OpConstantComposite %105 %28 %28
%116 = OpConstant %115 1 %107 = OpConstantComposite %97 %28 %28 %28
%117 = OpConstant %115 2 %108 = OpConstantComposite %97 %30 %30 %30
%118 = OpTypeVector %6 2 %109 = OpConstantComposite %7 %28 %28 %28 %28
%119 = OpConstantComposite %118 %21 %21 %110 = OpConstantComposite %7 %30 %30 %30 %30
%120 = OpConstantComposite %118 %62 %62 %122 = OpConstant %16 1
%121 = OpTypeVector %115 3 %123 = OpConstant %16 2
%122 = OpConstantComposite %121 %117 %117 %117 %124 = OpTypeVector %6 2
%123 = OpConstantComposite %121 %116 %116 %116 %125 = OpConstantComposite %124 %23 %23
%124 = OpConstantComposite %3 %53 %53 %53 %53 %126 = OpConstantComposite %124 %64 %64
%125 = OpConstantComposite %3 %15 %15 %15 %15 %127 = OpConstantComposite %15 %123 %123 %123
%126 = OpTypeVector %115 2 %128 = OpConstantComposite %15 %122 %122 %122
%127 = OpConstantComposite %126 %117 %117 %129 = OpConstantComposite %3 %58 %58 %58 %58
%128 = OpConstantComposite %126 %116 %116 %130 = OpConstantComposite %3 %17 %17 %17 %17
%129 = OpConstantNull %11 %131 = OpTypeVector %16 2
%130 = OpConstantNull %12 %132 = OpConstantComposite %131 %123 %123
%131 = OpConstantComposite %10 %53 %53 %53 %133 = OpConstantComposite %131 %122 %122
%132 = OpConstantNull %13 %134 = OpConstantNull %11
%296 = OpConstantNull %14 %135 = OpConstantNull %12
%298 = OpTypePointer Function %6 %136 = OpConstantComposite %10 %58 %58 %58
%299 = OpConstantNull %6 %137 = OpConstantNull %13
%301 = OpTypePointer Function %14 %301 = OpConstantNull %14
%329 = OpTypePointer Function %6 %303 = OpTypePointer Function %6
%367 = OpConstantComposite %10 %15 %15 %15 %304 = OpConstantNull %6
%24 = OpFunction %3 None %25 %306 = OpTypePointer Function %14
%23 = OpLabel %334 = OpTypePointer Function %6
OpBranch %32 %372 = OpTypePointer Input %15
%32 = OpLabel %371 = OpVariable %372 Input
%33 = OpSelect %6 %26 %21 %27 %375 = OpConstantComposite %10 %17 %17 %17
%35 = OpCompositeConstruct %7 %26 %26 %26 %26 %26 = OpFunction %3 None %27
%34 = OpSelect %3 %35 %16 %18 %25 = OpLabel
%36 = OpSelect %3 %29 %18 %16 OpBranch %34
%37 = OpExtInst %3 %1 FMix %18 %16 %20 %34 = OpLabel
%39 = OpCompositeConstruct %3 %30 %30 %30 %30 %35 = OpSelect %6 %28 %23 %29
%38 = OpExtInst %3 %1 FMix %18 %16 %39 %37 = OpCompositeConstruct %7 %28 %28 %28 %28
%40 = OpBitcast %4 %21 %36 = OpSelect %3 %37 %18 %20
%41 = OpBitcast %3 %22 %38 = OpSelect %3 %31 %20 %18
%42 = OpCompositeConstruct %5 %33 %33 %33 %33 %39 = OpExtInst %3 %1 FMix %20 %18 %22
%43 = OpIAdd %5 %42 %31 %41 = OpCompositeConstruct %3 %32 %32 %32 %32
%44 = OpConvertSToF %3 %43 %40 = OpExtInst %3 %1 FMix %20 %18 %41
%45 = OpFAdd %3 %44 %34 %42 = OpBitcast %4 %23
%46 = OpFAdd %3 %45 %37 %43 = OpBitcast %3 %24
%47 = OpFAdd %3 %46 %38 %44 = OpCompositeConstruct %5 %35 %35 %35 %35
%48 = OpCompositeConstruct %3 %40 %40 %40 %40 %45 = OpIAdd %5 %44 %33
%49 = OpFAdd %3 %47 %48 %46 = OpConvertSToF %3 %45
%50 = OpFAdd %3 %49 %41 %47 = OpFAdd %3 %46 %36
OpReturnValue %50 %48 = OpFAdd %3 %47 %39
%49 = OpFAdd %3 %48 %40
%50 = OpCompositeConstruct %3 %42 %42 %42 %42
%51 = OpFAdd %3 %49 %50
%52 = OpFAdd %3 %51 %43
OpReturnValue %52
OpFunctionEnd OpFunctionEnd
%52 = OpFunction %3 None %25 %56 = OpFunction %3 None %57
%51 = OpLabel %54 = OpFunctionParameter %4
OpBranch %64 %55 = OpFunctionParameter %6
%64 = OpLabel %53 = OpLabel
%65 = OpFAdd %9 %55 %54 OpBranch %66
%66 = OpFSub %9 %65 %57 %66 = OpLabel
%67 = OpFDiv %9 %66 %59 %67 = OpCompositeConstruct %9 %54 %54
%68 = OpSRem %5 %61 %63 %68 = OpFAdd %9 %59 %67
%69 = OpVectorShuffle %3 %67 %67 0 1 0 1 %69 = OpFSub %9 %68 %61
%70 = OpConvertSToF %3 %68 %70 = OpFDiv %9 %69 %63
%71 = OpFAdd %3 %69 %70 %71 = OpCompositeConstruct %5 %55 %55 %55 %55
OpReturnValue %71 %72 = OpSRem %5 %71 %65
%73 = OpVectorShuffle %3 %70 %70 0 1 0 1
%74 = OpConvertSToF %3 %72
%75 = OpFAdd %3 %73 %74
OpReturnValue %75
OpFunctionEnd OpFunctionEnd
%73 = OpFunction %9 None %74 %77 = OpFunction %9 None %78
%72 = OpLabel %76 = OpLabel
%75 = OpVariable %76 Function %54 %82 = OpVariable %83 Function %59
OpBranch %77 OpBranch %84
%77 = OpLabel %84 = OpLabel
%78 = OpLoad %9 %75 %85 = OpLoad %9 %82
%79 = OpFAdd %9 %78 %55 %86 = OpFAdd %9 %85 %79
OpStore %75 %79 OpStore %82 %86
%80 = OpLoad %9 %75 %87 = OpLoad %9 %82
%81 = OpFSub %9 %80 %57 %88 = OpFSub %9 %87 %81
OpStore %75 %81 OpStore %82 %88
%82 = OpLoad %9 %75 %89 = OpLoad %9 %82
%83 = OpFDiv %9 %82 %59 %90 = OpFDiv %9 %89 %61
OpStore %75 %83 OpStore %82 %90
%84 = OpLoad %9 %75 %91 = OpLoad %9 %82
OpReturnValue %84 OpReturnValue %91
OpFunctionEnd OpFunctionEnd
%87 = OpFunction %10 None %88 %94 = OpFunction %10 None %95
%86 = OpFunctionParameter %10 %93 = OpFunctionParameter %10
%85 = OpLabel %92 = OpLabel
OpBranch %89 OpBranch %96
%89 = OpLabel %96 = OpLabel
%92 = OpFUnordNotEqual %90 %86 %91 %99 = OpFUnordNotEqual %97 %93 %98
%94 = OpSelect %10 %92 %93 %91 %101 = OpSelect %10 %99 %100 %98
OpReturnValue %94 OpReturnValue %101
OpFunctionEnd OpFunctionEnd
%96 = OpFunction %2 None %97 %103 = OpFunction %2 None %104
%95 = OpLabel %102 = OpLabel
OpBranch %104 OpBranch %111
%104 = OpLabel %111 = OpLabel
%105 = OpLogicalNot %8 %26 %112 = OpLogicalNot %8 %28
%106 = OpLogicalNot %98 %99 %113 = OpLogicalNot %105 %106
%107 = OpLogicalOr %8 %26 %28 %114 = OpLogicalOr %8 %28 %30
%108 = OpLogicalAnd %8 %26 %28 %115 = OpLogicalAnd %8 %28 %30
%109 = OpLogicalOr %8 %26 %28 %116 = OpLogicalOr %8 %28 %30
%110 = OpLogicalOr %90 %100 %101 %117 = OpLogicalOr %97 %107 %108
%111 = OpLogicalAnd %8 %26 %28 %118 = OpLogicalAnd %8 %28 %30
%112 = OpLogicalAnd %7 %102 %103 %119 = OpLogicalAnd %7 %109 %110
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%114 = OpFunction %2 None %97 %121 = OpFunction %2 None %104
%113 = OpLabel %120 = OpLabel
OpBranch %133 OpBranch %138
%133 = OpLabel %138 = OpLabel
%134 = OpFNegate %4 %15 %139 = OpFNegate %4 %17
%135 = OpSNegate %118 %119 %140 = OpSNegate %124 %125
%136 = OpFNegate %9 %55 %141 = OpFNegate %9 %79
%137 = OpIAdd %6 %62 %21 %142 = OpIAdd %6 %64 %23
%138 = OpIAdd %115 %117 %116 %143 = OpIAdd %16 %123 %122
%139 = OpFAdd %4 %53 %15 %144 = OpFAdd %4 %58 %17
%140 = OpIAdd %118 %120 %119 %145 = OpIAdd %124 %126 %125
%141 = OpIAdd %121 %122 %123 %146 = OpIAdd %15 %127 %128
%142 = OpFAdd %3 %124 %125 %147 = OpFAdd %3 %129 %130
%143 = OpISub %6 %62 %21 %148 = OpISub %6 %64 %23
%144 = OpISub %115 %117 %116 %149 = OpISub %16 %123 %122
%145 = OpFSub %4 %53 %15 %150 = OpFSub %4 %58 %17
%146 = OpISub %118 %120 %119 %151 = OpISub %124 %126 %125
%147 = OpISub %121 %122 %123 %152 = OpISub %15 %127 %128
%148 = OpFSub %3 %124 %125 %153 = OpFSub %3 %129 %130
%149 = OpIMul %6 %62 %21 %154 = OpIMul %6 %64 %23
%150 = OpIMul %115 %117 %116 %155 = OpIMul %16 %123 %122
%151 = OpFMul %4 %53 %15 %156 = OpFMul %4 %58 %17
%152 = OpIMul %118 %120 %119 %157 = OpIMul %124 %126 %125
%153 = OpIMul %121 %122 %123 %158 = OpIMul %15 %127 %128
%154 = OpFMul %3 %124 %125 %159 = OpFMul %3 %129 %130
%155 = OpSDiv %6 %62 %21 %160 = OpSDiv %6 %64 %23
%156 = OpUDiv %115 %117 %116 %161 = OpUDiv %16 %123 %122
%157 = OpFDiv %4 %53 %15 %162 = OpFDiv %4 %58 %17
%158 = OpSDiv %118 %120 %119 %163 = OpSDiv %124 %126 %125
%159 = OpUDiv %121 %122 %123 %164 = OpUDiv %15 %127 %128
%160 = OpFDiv %3 %124 %125 %165 = OpFDiv %3 %129 %130
%161 = OpSRem %6 %62 %21 %166 = OpSRem %6 %64 %23
%162 = OpUMod %115 %117 %116 %167 = OpUMod %16 %123 %122
%163 = OpFRem %4 %53 %15 %168 = OpFRem %4 %58 %17
%164 = OpSRem %118 %120 %119 %169 = OpSRem %124 %126 %125
%165 = OpUMod %121 %122 %123 %170 = OpUMod %15 %127 %128
%166 = OpFRem %3 %124 %125 %171 = OpFRem %3 %129 %130
OpBranch %167 OpBranch %172
%167 = OpLabel %172 = OpLabel
%169 = OpIAdd %118 %120 %119 %174 = OpIAdd %124 %126 %125
%170 = OpIAdd %118 %120 %119 %175 = OpIAdd %124 %126 %125
%171 = OpIAdd %126 %127 %128 %176 = OpIAdd %131 %132 %133
%172 = OpIAdd %126 %127 %128 %177 = OpIAdd %131 %132 %133
%173 = OpFAdd %9 %54 %55 %178 = OpFAdd %9 %59 %79
%174 = OpFAdd %9 %54 %55 %179 = OpFAdd %9 %59 %79
%175 = OpISub %118 %120 %119 %180 = OpISub %124 %126 %125
%176 = OpISub %118 %120 %119 %181 = OpISub %124 %126 %125
%177 = OpISub %126 %127 %128 %182 = OpISub %131 %132 %133
%178 = OpISub %126 %127 %128 %183 = OpISub %131 %132 %133
%179 = OpFSub %9 %54 %55 %184 = OpFSub %9 %59 %79
%180 = OpFSub %9 %54 %55 %185 = OpFSub %9 %59 %79
%182 = OpCompositeConstruct %118 %21 %21 %187 = OpCompositeConstruct %124 %23 %23
%181 = OpIMul %118 %120 %182 %186 = OpIMul %124 %126 %187
%184 = OpCompositeConstruct %118 %62 %62 %189 = OpCompositeConstruct %124 %64 %64
%183 = OpIMul %118 %119 %184 %188 = OpIMul %124 %125 %189
%186 = OpCompositeConstruct %126 %116 %116 %191 = OpCompositeConstruct %131 %122 %122
%185 = OpIMul %126 %127 %186 %190 = OpIMul %131 %132 %191
%188 = OpCompositeConstruct %126 %117 %117 %193 = OpCompositeConstruct %131 %123 %123
%187 = OpIMul %126 %128 %188 %192 = OpIMul %131 %133 %193
%189 = OpVectorTimesScalar %9 %54 %15 %194 = OpVectorTimesScalar %9 %59 %17
%190 = OpVectorTimesScalar %9 %55 %53 %195 = OpVectorTimesScalar %9 %79 %58
%191 = OpSDiv %118 %120 %119 %196 = OpSDiv %124 %126 %125
%192 = OpSDiv %118 %120 %119 %197 = OpSDiv %124 %126 %125
%193 = OpUDiv %126 %127 %128 %198 = OpUDiv %131 %132 %133
%194 = OpUDiv %126 %127 %128 %199 = OpUDiv %131 %132 %133
%195 = OpFDiv %9 %54 %55 %200 = OpFDiv %9 %59 %79
%196 = OpFDiv %9 %54 %55 %201 = OpFDiv %9 %59 %79
%197 = OpSRem %118 %120 %119 %202 = OpSRem %124 %126 %125
%198 = OpSRem %118 %120 %119 %203 = OpSRem %124 %126 %125
%199 = OpUMod %126 %127 %128 %204 = OpUMod %131 %132 %133
%200 = OpUMod %126 %127 %128 %205 = OpUMod %131 %132 %133
%201 = OpFRem %9 %54 %55 %206 = OpFRem %9 %59 %79
%202 = OpFRem %9 %54 %55 %207 = OpFRem %9 %59 %79
OpBranch %168 OpBranch %173
%168 = OpLabel %173 = OpLabel
%204 = OpCompositeExtract %10 %129 0 %209 = OpCompositeExtract %10 %134 0
%205 = OpCompositeExtract %10 %129 0 %210 = OpCompositeExtract %10 %134 0
%206 = OpFAdd %10 %204 %205 %211 = OpFAdd %10 %209 %210
%207 = OpCompositeExtract %10 %129 1 %212 = OpCompositeExtract %10 %134 1
%208 = OpCompositeExtract %10 %129 1 %213 = OpCompositeExtract %10 %134 1
%209 = OpFAdd %10 %207 %208 %214 = OpFAdd %10 %212 %213
%210 = OpCompositeExtract %10 %129 2 %215 = OpCompositeExtract %10 %134 2
%211 = OpCompositeExtract %10 %129 2 %216 = OpCompositeExtract %10 %134 2
%212 = OpFAdd %10 %210 %211 %217 = OpFAdd %10 %215 %216
%203 = OpCompositeConstruct %11 %206 %209 %212 %208 = OpCompositeConstruct %11 %211 %214 %217
%214 = OpCompositeExtract %10 %129 0 %219 = OpCompositeExtract %10 %134 0
%215 = OpCompositeExtract %10 %129 0 %220 = OpCompositeExtract %10 %134 0
%216 = OpFSub %10 %214 %215 %221 = OpFSub %10 %219 %220
%217 = OpCompositeExtract %10 %129 1 %222 = OpCompositeExtract %10 %134 1
%218 = OpCompositeExtract %10 %129 1 %223 = OpCompositeExtract %10 %134 1
%219 = OpFSub %10 %217 %218 %224 = OpFSub %10 %222 %223
%220 = OpCompositeExtract %10 %129 2 %225 = OpCompositeExtract %10 %134 2
%221 = OpCompositeExtract %10 %129 2 %226 = OpCompositeExtract %10 %134 2
%222 = OpFSub %10 %220 %221 %227 = OpFSub %10 %225 %226
%213 = OpCompositeConstruct %11 %216 %219 %222 %218 = OpCompositeConstruct %11 %221 %224 %227
%223 = OpMatrixTimesScalar %11 %129 %15 %228 = OpMatrixTimesScalar %11 %134 %17
%224 = OpMatrixTimesScalar %11 %129 %53 %229 = OpMatrixTimesScalar %11 %134 %58
%225 = OpMatrixTimesVector %10 %130 %125 %230 = OpMatrixTimesVector %10 %135 %130
%226 = OpVectorTimesMatrix %3 %131 %130 %231 = OpVectorTimesMatrix %3 %136 %135
%227 = OpMatrixTimesMatrix %11 %130 %132 %232 = OpMatrixTimesMatrix %11 %135 %137
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%229 = OpFunction %2 None %97 %234 = OpFunction %2 None %104
%228 = OpLabel %233 = OpLabel
OpBranch %230 OpBranch %235
%230 = OpLabel %235 = OpLabel
%231 = OpNot %6 %21 %236 = OpNot %6 %23
%232 = OpNot %115 %116 %237 = OpNot %16 %122
%233 = OpNot %118 %119 %238 = OpNot %124 %125
%234 = OpNot %121 %123 %239 = OpNot %15 %128
%235 = OpBitwiseOr %6 %62 %21 %240 = OpBitwiseOr %6 %64 %23
%236 = OpBitwiseOr %115 %117 %116 %241 = OpBitwiseOr %16 %123 %122
%237 = OpBitwiseOr %118 %120 %119 %242 = OpBitwiseOr %124 %126 %125
%238 = OpBitwiseOr %121 %122 %123 %243 = OpBitwiseOr %15 %127 %128
%239 = OpBitwiseAnd %6 %62 %21 %244 = OpBitwiseAnd %6 %64 %23
%240 = OpBitwiseAnd %115 %117 %116 %245 = OpBitwiseAnd %16 %123 %122
%241 = OpBitwiseAnd %118 %120 %119 %246 = OpBitwiseAnd %124 %126 %125
%242 = OpBitwiseAnd %121 %122 %123 %247 = OpBitwiseAnd %15 %127 %128
%243 = OpBitwiseXor %6 %62 %21 %248 = OpBitwiseXor %6 %64 %23
%244 = OpBitwiseXor %115 %117 %116 %249 = OpBitwiseXor %16 %123 %122
%245 = OpBitwiseXor %118 %120 %119 %250 = OpBitwiseXor %124 %126 %125
%246 = OpBitwiseXor %121 %122 %123 %251 = OpBitwiseXor %15 %127 %128
%247 = OpShiftLeftLogical %6 %62 %116 %252 = OpShiftLeftLogical %6 %64 %122
%248 = OpShiftLeftLogical %115 %117 %116 %253 = OpShiftLeftLogical %16 %123 %122
%249 = OpShiftLeftLogical %118 %120 %128 %254 = OpShiftLeftLogical %124 %126 %133
%250 = OpShiftLeftLogical %121 %122 %123 %255 = OpShiftLeftLogical %15 %127 %128
%251 = OpShiftRightArithmetic %6 %62 %116 %256 = OpShiftRightArithmetic %6 %64 %122
%252 = OpShiftRightLogical %115 %117 %116 %257 = OpShiftRightLogical %16 %123 %122
%253 = OpShiftRightArithmetic %118 %120 %128 %258 = OpShiftRightArithmetic %124 %126 %133
%254 = OpShiftRightLogical %121 %122 %123 %259 = OpShiftRightLogical %15 %127 %128
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%256 = OpFunction %2 None %97 %261 = OpFunction %2 None %104
%255 = OpLabel %260 = OpLabel
OpBranch %257 OpBranch %262
%257 = OpLabel %262 = OpLabel
%258 = OpIEqual %8 %62 %21 %263 = OpIEqual %8 %64 %23
%259 = OpIEqual %8 %117 %116 %264 = OpIEqual %8 %123 %122
%260 = OpFOrdEqual %8 %53 %15 %265 = OpFOrdEqual %8 %58 %17
%261 = OpIEqual %98 %120 %119 %266 = OpIEqual %105 %126 %125
%262 = OpIEqual %90 %122 %123 %267 = OpIEqual %97 %127 %128
%263 = OpFOrdEqual %7 %124 %125 %268 = OpFOrdEqual %7 %129 %130
%264 = OpINotEqual %8 %62 %21 %269 = OpINotEqual %8 %64 %23
%265 = OpINotEqual %8 %117 %116 %270 = OpINotEqual %8 %123 %122
%266 = OpFOrdNotEqual %8 %53 %15 %271 = OpFOrdNotEqual %8 %58 %17
%267 = OpINotEqual %98 %120 %119 %272 = OpINotEqual %105 %126 %125
%268 = OpINotEqual %90 %122 %123 %273 = OpINotEqual %97 %127 %128
%269 = OpFOrdNotEqual %7 %124 %125 %274 = OpFOrdNotEqual %7 %129 %130
%270 = OpSLessThan %8 %62 %21 %275 = OpSLessThan %8 %64 %23
%271 = OpULessThan %8 %117 %116 %276 = OpULessThan %8 %123 %122
%272 = OpFOrdLessThan %8 %53 %15 %277 = OpFOrdLessThan %8 %58 %17
%273 = OpSLessThan %98 %120 %119 %278 = OpSLessThan %105 %126 %125
%274 = OpULessThan %90 %122 %123 %279 = OpULessThan %97 %127 %128
%275 = OpFOrdLessThan %7 %124 %125 %280 = OpFOrdLessThan %7 %129 %130
%276 = OpSLessThanEqual %8 %62 %21 %281 = OpSLessThanEqual %8 %64 %23
%277 = OpULessThanEqual %8 %117 %116 %282 = OpULessThanEqual %8 %123 %122
%278 = OpFOrdLessThanEqual %8 %53 %15 %283 = OpFOrdLessThanEqual %8 %58 %17
%279 = OpSLessThanEqual %98 %120 %119 %284 = OpSLessThanEqual %105 %126 %125
%280 = OpULessThanEqual %90 %122 %123 %285 = OpULessThanEqual %97 %127 %128
%281 = OpFOrdLessThanEqual %7 %124 %125 %286 = OpFOrdLessThanEqual %7 %129 %130
%282 = OpSGreaterThan %8 %62 %21 %287 = OpSGreaterThan %8 %64 %23
%283 = OpUGreaterThan %8 %117 %116 %288 = OpUGreaterThan %8 %123 %122
%284 = OpFOrdGreaterThan %8 %53 %15 %289 = OpFOrdGreaterThan %8 %58 %17
%285 = OpSGreaterThan %98 %120 %119 %290 = OpSGreaterThan %105 %126 %125
%286 = OpUGreaterThan %90 %122 %123 %291 = OpUGreaterThan %97 %127 %128
%287 = OpFOrdGreaterThan %7 %124 %125 %292 = OpFOrdGreaterThan %7 %129 %130
%288 = OpSGreaterThanEqual %8 %62 %21 %293 = OpSGreaterThanEqual %8 %64 %23
%289 = OpUGreaterThanEqual %8 %117 %116 %294 = OpUGreaterThanEqual %8 %123 %122
%290 = OpFOrdGreaterThanEqual %8 %53 %15 %295 = OpFOrdGreaterThanEqual %8 %58 %17
%291 = OpSGreaterThanEqual %98 %120 %119 %296 = OpSGreaterThanEqual %105 %126 %125
%292 = OpUGreaterThanEqual %90 %122 %123 %297 = OpUGreaterThanEqual %97 %127 %128
%293 = OpFOrdGreaterThanEqual %7 %124 %125 %298 = OpFOrdGreaterThanEqual %7 %129 %130
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%295 = OpFunction %2 None %97 %300 = OpFunction %2 None %104
%294 = OpLabel %299 = OpLabel
%297 = OpVariable %298 Function %299 %302 = OpVariable %303 Function %304
%300 = OpVariable %301 Function %296 %305 = OpVariable %306 Function %301
OpBranch %302 OpBranch %307
%302 = OpLabel %307 = OpLabel
OpStore %297 %21 OpStore %302 %23
%303 = OpLoad %6 %297 %308 = OpLoad %6 %302
%304 = OpIAdd %6 %303 %21 %309 = OpIAdd %6 %308 %23
OpStore %297 %304 OpStore %302 %309
%305 = OpLoad %6 %297 %310 = OpLoad %6 %302
%306 = OpISub %6 %305 %21 %311 = OpISub %6 %310 %23
OpStore %297 %306 OpStore %302 %311
%307 = OpLoad %6 %297 %312 = OpLoad %6 %302
%308 = OpLoad %6 %297 %313 = OpLoad %6 %302
%309 = OpIMul %6 %308 %307 %314 = OpIMul %6 %313 %312
OpStore %297 %309 OpStore %302 %314
%310 = OpLoad %6 %297 %315 = OpLoad %6 %302
%311 = OpLoad %6 %297 %316 = OpLoad %6 %302
%312 = OpSDiv %6 %311 %310 %317 = OpSDiv %6 %316 %315
OpStore %297 %312 OpStore %302 %317
%313 = OpLoad %6 %297 %318 = OpLoad %6 %302
%314 = OpSRem %6 %313 %21 %319 = OpSRem %6 %318 %23
OpStore %297 %314 OpStore %302 %319
%315 = OpLoad %6 %297 %320 = OpLoad %6 %302
%316 = OpBitwiseAnd %6 %315 %27 %321 = OpBitwiseAnd %6 %320 %29
OpStore %297 %316 OpStore %302 %321
%317 = OpLoad %6 %297 %322 = OpLoad %6 %302
%318 = OpBitwiseOr %6 %317 %27 %323 = OpBitwiseOr %6 %322 %29
OpStore %297 %318 OpStore %302 %323
%319 = OpLoad %6 %297 %324 = OpLoad %6 %302
%320 = OpBitwiseXor %6 %319 %27 %325 = OpBitwiseXor %6 %324 %29
OpStore %297 %320 OpStore %302 %325
%321 = OpLoad %6 %297 %326 = OpLoad %6 %302
%322 = OpShiftLeftLogical %6 %321 %117 %327 = OpShiftLeftLogical %6 %326 %123
OpStore %297 %322 OpStore %302 %327
%323 = OpLoad %6 %297 %328 = OpLoad %6 %302
%324 = OpShiftRightArithmetic %6 %323 %116 %329 = OpShiftRightArithmetic %6 %328 %122
OpStore %297 %324 OpStore %302 %329
%325 = OpLoad %6 %297 %330 = OpLoad %6 %302
%326 = OpIAdd %6 %325 %21 %331 = OpIAdd %6 %330 %23
OpStore %297 %326 OpStore %302 %331
%327 = OpLoad %6 %297 %332 = OpLoad %6 %302
%328 = OpISub %6 %327 %21 %333 = OpISub %6 %332 %23
OpStore %297 %328 OpStore %302 %333
%330 = OpAccessChain %329 %300 %21 %335 = OpAccessChain %334 %305 %23
%331 = OpLoad %6 %330 %336 = OpLoad %6 %335
%332 = OpIAdd %6 %331 %21 %337 = OpIAdd %6 %336 %23
%333 = OpAccessChain %329 %300 %21 %338 = OpAccessChain %334 %305 %23
OpStore %333 %332 OpStore %338 %337
%334 = OpAccessChain %329 %300 %21 %339 = OpAccessChain %334 %305 %23
%335 = OpLoad %6 %334 %340 = OpLoad %6 %339
%336 = OpISub %6 %335 %21 %341 = OpISub %6 %340 %23
%337 = OpAccessChain %329 %300 %21 %342 = OpAccessChain %334 %305 %23
OpStore %337 %336 OpStore %342 %341
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%339 = OpFunction %2 None %97 %344 = OpFunction %2 None %104
%338 = OpLabel %343 = OpLabel
OpBranch %340 OpBranch %345
%340 = OpLabel %345 = OpLabel
%341 = OpSNegate %6 %21 %346 = OpSNegate %6 %23
%342 = OpSNegate %6 %21 %347 = OpSNegate %6 %23
%343 = OpSNegate %6 %342 %348 = OpSNegate %6 %347
%344 = OpSNegate %6 %21 %349 = OpSNegate %6 %23
%345 = OpSNegate %6 %344
%346 = OpSNegate %6 %21
%347 = OpSNegate %6 %346
%348 = OpSNegate %6 %21
%349 = OpSNegate %6 %348
%350 = OpSNegate %6 %349 %350 = OpSNegate %6 %349
%351 = OpSNegate %6 %21 %351 = OpSNegate %6 %23
%352 = OpSNegate %6 %351 %352 = OpSNegate %6 %351
%353 = OpSNegate %6 %352 %353 = OpSNegate %6 %23
%354 = OpSNegate %6 %353 %354 = OpSNegate %6 %353
%355 = OpSNegate %6 %21 %355 = OpSNegate %6 %354
%356 = OpSNegate %6 %355 %356 = OpSNegate %6 %23
%357 = OpSNegate %6 %356 %357 = OpSNegate %6 %356
%358 = OpSNegate %6 %357 %358 = OpSNegate %6 %357
%359 = OpSNegate %6 %358 %359 = OpSNegate %6 %358
%360 = OpSNegate %6 %21 %360 = OpSNegate %6 %23
%361 = OpSNegate %6 %360 %361 = OpSNegate %6 %360
%362 = OpSNegate %6 %361 %362 = OpSNegate %6 %361
%363 = OpSNegate %6 %362 %363 = OpSNegate %6 %362
%364 = OpSNegate %6 %363 %364 = OpSNegate %6 %363
%365 = OpSNegate %6 %23
%366 = OpSNegate %6 %365
%367 = OpSNegate %6 %366
%368 = OpSNegate %6 %367
%369 = OpSNegate %6 %368
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%366 = OpFunction %2 None %97 %374 = OpFunction %2 None %104
%365 = OpLabel %370 = OpLabel
OpBranch %368 %373 = OpLoad %15 %371
%368 = OpLabel OpBranch %376
%369 = OpFunctionCall %3 %24 %376 = OpLabel
%370 = OpFunctionCall %3 %52 %377 = OpFunctionCall %3 %26
%371 = OpFunctionCall %10 %87 %367 %378 = OpCompositeExtract %16 %373 0
%372 = OpFunctionCall %2 %96 %379 = OpConvertUToF %4 %378
%373 = OpFunctionCall %2 %114 %380 = OpCompositeExtract %16 %373 1
%374 = OpFunctionCall %2 %229 %381 = OpBitcast %6 %380
%375 = OpFunctionCall %2 %256 %382 = OpFunctionCall %3 %56 %379 %381
%376 = OpFunctionCall %2 %295 %383 = OpFunctionCall %10 %94 %375
%384 = OpFunctionCall %2 %103
%385 = OpFunctionCall %2 %121
%386 = OpFunctionCall %2 %234
%387 = OpFunctionCall %2 %261
%388 = OpFunctionCall %2 %300
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

View File

@ -15,9 +15,9 @@ fn builtins() -> vec4<f32> {
return (((((vec4<f32>((vec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_); return (((((vec4<f32>((vec4(s1_) + v_i32_zero)) + s2_) + m1_) + m2_) + vec4(b1_)) + b2_);
} }
fn splat() -> vec4<f32> { fn splat(m: f32, n: i32) -> vec4<f32> {
let a_2 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0)); let a_2 = (((vec2(2.0) + vec2(m)) - vec2(4.0)) / vec2(8.0));
let b = (vec4(5) % vec4(2)); let b = (vec4(n) % vec4(2));
return (a_2.xyxy + vec4<f32>(b)); return (a_2.xyxy + vec4<f32>(b));
} }
@ -241,10 +241,10 @@ fn negation_avoids_prefix_decrement() {
} }
@compute @workgroup_size(1, 1, 1) @compute @workgroup_size(1, 1, 1)
fn main() { fn main(@builtin(workgroup_id) id: vec3<u32>) {
let _e0 = builtins(); let _e1 = builtins();
let _e1 = splat(); let _e6 = splat(f32(id.x), i32(id.y));
let _e6 = bool_cast(vec3<f32>(1.0, 1.0, 1.0)); let _e11 = bool_cast(vec3<f32>(1.0, 1.0, 1.0));
logical(); logical();
arithmetic(); arithmetic();
bit(); bit();