Add snapshot tests for constant evaluation of splats and composes.

This commit is contained in:
Jim Blandy
2023-09-28 16:31:59 -07:00
committed by Teodor Tanasoaia
parent c33d7ee40d
commit 6bd15e4492
6 changed files with 143 additions and 46 deletions

View File

@@ -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);
}

View File

@@ -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;
}

View File

@@ -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;
}

View File

@@ -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;
}

View File

@@ -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

View File

@@ -1,3 +1,5 @@
const FOUR: i32 = 4;
@group(0) @binding(0)
var<storage, read_write> out: vec4<i32>;
@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<i32>(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;
}