Split const-exprs.rs test into separate functions.

Also, just use a single out variable for each type. rather than
introducing a new output variable for every test we add.
This commit is contained in:
Jim Blandy 2023-09-20 16:01:00 -07:00 committed by Teodor Tanasoaia
parent 1a4b526d9a
commit bdcb9f6f64
6 changed files with 174 additions and 61 deletions

View File

@ -1,14 +1,28 @@
@group(0) @binding(0) var<storage, read_write> out: vec4<i32>;
@group(0) @binding(1) var<storage, read_write> out2: i32;
@group(0) @binding(2) var<storage, read_write> out3: i32;
@compute @workgroup_size(1)
fn main() {
swizzle_of_compose();
index_of_compose();
compose_three_deep();
}
// Swizzle the value of nested Compose expressions.
fn swizzle_of_compose() {
let a = vec2(1, 2);
let b = vec2(3, 4);
out = vec4(a, b).wzyx;
out2 = vec4(a, b)[1];
out3 = vec4(vec3(vec2(6, 7), 8), 9)[0];
out = vec4(a, b).wzyx; // should assign vec4(4, 3, 2, 1);
}
// Index the value of nested Compose expressions.
fn index_of_compose() {
let a = vec2(1, 2);
let b = vec2(3, 4);
out2 += vec4(a, b)[1]; // should assign 2
}
// Index the value of Compose expressions nested three deep
fn compose_three_deep() {
out2 += vec4(vec3(vec2(6, 7), 8), 9)[0]; // should assign 6
}

View File

@ -9,15 +9,32 @@ 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; };
layout(std430) buffer type_1_block_2Compute { int _group_0_binding_2_cs; };
void main() {
void swizzle_of_compose() {
ivec2 a = ivec2(1, 2);
ivec2 b = ivec2(3, 4);
_group_0_binding_0_cs = ivec4(4, 3, 2, 1);
_group_0_binding_1_cs = 2;
_group_0_binding_2_cs = 6;
return;
}
void index_of_compose() {
ivec2 a_1 = ivec2(1, 2);
ivec2 b_1 = ivec2(3, 4);
int _e7 = _group_0_binding_1_cs;
_group_0_binding_1_cs = (_e7 + 2);
return;
}
void compose_three_deep() {
int _e2 = _group_0_binding_1_cs;
_group_0_binding_1_cs = (_e2 + 6);
return;
}
void main() {
swizzle_of_compose();
index_of_compose();
compose_three_deep();
return;
}

View File

@ -1,14 +1,35 @@
RWByteAddressBuffer out_ : register(u0);
RWByteAddressBuffer out2_ : register(u1);
RWByteAddressBuffer out3_ : register(u2);
[numthreads(1, 1, 1)]
void main()
void swizzle_of_compose()
{
int2 a = int2(1, 2);
int2 b = int2(3, 4);
out_.Store4(0, asuint(int4(4, 3, 2, 1)));
out2_.Store(0, asuint(2));
out3_.Store(0, asuint(6));
return;
}
void index_of_compose()
{
int2 a_1 = int2(1, 2);
int2 b_1 = int2(3, 4);
int _expr7 = asint(out2_.Load(0));
out2_.Store(0, asuint((_expr7 + 2)));
return;
}
void compose_three_deep()
{
int _expr2 = asint(out2_.Load(0));
out2_.Store(0, asuint((_expr2 + 6)));
return;
}
[numthreads(1, 1, 1)]
void main()
{
swizzle_of_compose();
index_of_compose();
compose_three_deep();
return;
}

View File

@ -5,15 +5,39 @@
using metal::uint;
kernel void main_(
device metal::int4& out [[user(fake0)]]
, device int& out2_ [[user(fake0)]]
, device int& out3_ [[user(fake0)]]
void swizzle_of_compose(
device metal::int4& out
) {
metal::int2 a = metal::int2(1, 2);
metal::int2 b = metal::int2(3, 4);
out = metal::int4(4, 3, 2, 1);
out2_ = 2;
out3_ = 6;
return;
}
void index_of_compose(
device int& out2_
) {
metal::int2 a_1 = metal::int2(1, 2);
metal::int2 b_1 = metal::int2(3, 4);
int _e7 = out2_;
out2_ = _e7 + 2;
return;
}
void compose_three_deep(
device int& out2_
) {
int _e2 = out2_;
out2_ = _e2 + 6;
return;
}
kernel void main_(
device metal::int4& out [[user(fake0)]]
, device int& out2_ [[user(fake0)]]
) {
swizzle_of_compose(out);
index_of_compose(out2_);
compose_three_deep(out2_);
return;
}

View File

@ -1,13 +1,13 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 34
; Bound: 49
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %16 "main"
OpExecutionMode %16 LocalSize 1 1 1
OpEntryPoint GLCompute %42 "main"
OpExecutionMode %42 LocalSize 1 1 1
OpDecorate %6 DescriptorSet 0
OpDecorate %6 Binding 0
OpDecorate %7 Block
@ -16,10 +16,6 @@ OpDecorate %9 DescriptorSet 0
OpDecorate %9 Binding 1
OpDecorate %10 Block
OpMemberDecorate %10 0 Offset 0
OpDecorate %12 DescriptorSet 0
OpDecorate %12 Binding 2
OpDecorate %13 Block
OpMemberDecorate %13 0 Offset 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpTypeVector %4 4
@ -30,31 +26,55 @@ OpMemberDecorate %13 0 Offset 0
%10 = OpTypeStruct %4
%11 = OpTypePointer StorageBuffer %10
%9 = OpVariable %11 StorageBuffer
%13 = OpTypeStruct %4
%14 = OpTypePointer StorageBuffer %13
%12 = OpVariable %14 StorageBuffer
%17 = OpTypeFunction %2
%18 = OpTypePointer StorageBuffer %3
%20 = OpTypeInt 32 0
%19 = OpConstant %20 0
%22 = OpTypePointer StorageBuffer %4
%25 = OpConstant %4 1
%26 = OpConstant %4 2
%27 = OpConstantComposite %5 %25 %26
%28 = OpConstant %4 3
%29 = OpConstant %4 4
%30 = OpConstantComposite %5 %28 %29
%31 = OpConstantComposite %3 %29 %28 %26 %25
%32 = OpConstant %4 6
%16 = OpFunction %2 None %17
%15 = OpLabel
%21 = OpAccessChain %18 %6 %19
%23 = OpAccessChain %22 %9 %19
%24 = OpAccessChain %22 %12 %19
OpBranch %33
%33 = OpLabel
OpStore %21 %31
OpStore %23 %26
OpStore %24 %32
%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
%29 = OpTypePointer StorageBuffer %4
%37 = OpConstant %4 6
%13 = OpFunction %2 None %14
%12 = OpLabel
%18 = OpAccessChain %15 %6 %16
OpBranch %26
%26 = OpLabel
OpStore %18 %25
OpReturn
OpFunctionEnd
%28 = OpFunction %2 None %14
%27 = OpLabel
%30 = OpAccessChain %29 %9 %16
OpBranch %31
%31 = OpLabel
%32 = OpLoad %4 %30
%33 = OpIAdd %4 %32 %20
OpStore %30 %33
OpReturn
OpFunctionEnd
%35 = OpFunction %2 None %14
%34 = OpLabel
%36 = OpAccessChain %29 %9 %16
OpBranch %38
%38 = OpLabel
%39 = OpLoad %4 %36
%40 = OpIAdd %4 %39 %37
OpStore %36 %40
OpReturn
OpFunctionEnd
%42 = OpFunction %2 None %14
%41 = OpLabel
%43 = OpAccessChain %15 %6 %16
%44 = OpAccessChain %29 %9 %16
OpBranch %45
%45 = OpLabel
%46 = OpFunctionCall %2 %13
%47 = OpFunctionCall %2 %28
%48 = OpFunctionCall %2 %35
OpReturn
OpFunctionEnd

View File

@ -2,15 +2,32 @@
var<storage, read_write> out: vec4<i32>;
@group(0) @binding(1)
var<storage, read_write> out2_: i32;
@group(0) @binding(2)
var<storage, read_write> out3_: i32;
@compute @workgroup_size(1, 1, 1)
fn main() {
fn swizzle_of_compose() {
let a = vec2<i32>(1, 2);
let b = vec2<i32>(3, 4);
out = vec4<i32>(4, 3, 2, 1);
out2_ = 2;
out3_ = 6;
return;
}
fn index_of_compose() {
let a_1 = vec2<i32>(1, 2);
let b_1 = vec2<i32>(3, 4);
let _e7 = out2_;
out2_ = (_e7 + 2);
return;
}
fn compose_three_deep() {
let _e2 = out2_;
out2_ = (_e2 + 6);
return;
}
@compute @workgroup_size(1, 1, 1)
fn main() {
swizzle_of_compose();
index_of_compose();
compose_three_deep();
return;
}