mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-12-10 22:17:51 +00:00
Add tests derived from VK-GL-CTS
This adds SPIR-V assembly and WGSL tests derived from VK-GL-CTS commit 571256871c2e2f03995373e1e4a02958d8cd8cf5. The following procedure was followed: - Those .amber files in VK-GL-CTS wholly owned by Google were identified - All GLSL and SPIR-V shaders were extracted from the Amber files and converted into SPIR-V binaries - The compact-ids pass of spirv-opt was applied to each binary - Duplicate binaries were removed - spirv-opt -O was used to obtain an optimized version of each remaining binary, with duplicates discarded - Binaries that failed validation using spirv-val with target environment SPIR-V 1.3 were discarded - Those binaries that tint could not successfully convert into WGSL were put aside for further investigation - SPIR-V assembly versions of the remaining binaries are included in this CL - test-runner with -generate-expected and -generate-skip was used to generate expected .spvasm, .msl, .hlsl and .wgsl outputs for these SPIR-V assembly tests - Each successfully-generated .expected.wgsl is included in this CL again, as a WGLSL test - test-runner with -generate-expected and -generate-skip was used again, to generate expected outputs for these WGSL tests Change-Id: Ibe9baf2729cf97e0b633db9a426f53362a5de540 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58842 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
This commit is contained in:
Alastair Donaldson
committed by
Tint LUCI CQ
Tint LUCI CQ
parent
bd3edb564f
commit
f7e73d4ee3
@@ -0,0 +1,209 @@
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_WorkGroupID
|
||||
OpExecutionMode %main LocalSize 4 1 1
|
||||
OpSource GLSL 430
|
||||
OpName %main "main"
|
||||
OpName %base_index_in "base_index_in"
|
||||
OpName %gl_WorkGroupID "gl_WorkGroupID"
|
||||
OpName %base_index_out "base_index_out"
|
||||
OpName %index_in0 "index_in0"
|
||||
OpName %index_in1 "index_in1"
|
||||
OpName %index_out0 "index_out0"
|
||||
OpName %index_out1 "index_out1"
|
||||
OpName %condition_index "condition_index"
|
||||
OpName %i "i"
|
||||
OpName %In2 "In2"
|
||||
OpMemberName %In2 0 "data_in2"
|
||||
OpName %_ ""
|
||||
OpName %Out0 "Out0"
|
||||
OpMemberName %Out0 0 "data_out0"
|
||||
OpName %__0 ""
|
||||
OpName %In1 "In1"
|
||||
OpMemberName %In1 0 "data_in1"
|
||||
OpName %__1 ""
|
||||
OpName %In0 "In0"
|
||||
OpMemberName %In0 0 "data_in0"
|
||||
OpName %__2 ""
|
||||
OpName %temp0 "temp0"
|
||||
OpName %temp1 "temp1"
|
||||
OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId
|
||||
OpDecorate %_arr_int_uint_8 ArrayStride 4
|
||||
OpMemberDecorate %In2 0 NonWritable
|
||||
OpMemberDecorate %In2 0 Offset 0
|
||||
OpDecorate %In2 Block
|
||||
OpDecorate %_ DescriptorSet 0
|
||||
OpDecorate %_ Binding 2
|
||||
OpDecorate %_arr_int_uint_1024 ArrayStride 4
|
||||
OpMemberDecorate %Out0 0 NonReadable
|
||||
OpMemberDecorate %Out0 0 Offset 0
|
||||
OpDecorate %Out0 Block
|
||||
OpDecorate %__0 DescriptorSet 0
|
||||
OpDecorate %__0 Binding 3
|
||||
OpDecorate %_arr_int_uint_512 ArrayStride 4
|
||||
OpMemberDecorate %In1 0 NonWritable
|
||||
OpMemberDecorate %In1 0 Offset 0
|
||||
OpDecorate %In1 Block
|
||||
OpDecorate %__1 DescriptorSet 0
|
||||
OpDecorate %__1 Binding 1
|
||||
OpDecorate %_arr_int_uint_512_0 ArrayStride 4
|
||||
OpMemberDecorate %In0 0 NonWritable
|
||||
OpMemberDecorate %In0 0 Offset 0
|
||||
OpDecorate %In0 Block
|
||||
OpDecorate %__2 DescriptorSet 0
|
||||
OpDecorate %__2 Binding 0
|
||||
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
|
||||
%void = OpTypeVoid
|
||||
%28 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%uint_128 = OpConstant %uint 128
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%uint_256 = OpConstant %uint 256
|
||||
%int = OpTypeInt 32 1
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%int_127 = OpConstant %int 127
|
||||
%int_383 = OpConstant %int 383
|
||||
%int_255 = OpConstant %int 255
|
||||
%int_0 = OpConstant %int 0
|
||||
%int_256 = OpConstant %int 256
|
||||
%bool = OpTypeBool
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_int_uint_8 = OpTypeArray %int %uint_8
|
||||
%In2 = OpTypeStruct %_arr_int_uint_8
|
||||
%_ptr_StorageBuffer_In2 = OpTypePointer StorageBuffer %In2
|
||||
%_ = OpVariable %_ptr_StorageBuffer_In2 StorageBuffer
|
||||
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
|
||||
%uint_1024 = OpConstant %uint 1024
|
||||
%_arr_int_uint_1024 = OpTypeArray %int %uint_1024
|
||||
%Out0 = OpTypeStruct %_arr_int_uint_1024
|
||||
%_ptr_StorageBuffer_Out0 = OpTypePointer StorageBuffer %Out0
|
||||
%__0 = OpVariable %_ptr_StorageBuffer_Out0 StorageBuffer
|
||||
%uint_512 = OpConstant %uint 512
|
||||
%_arr_int_uint_512 = OpTypeArray %int %uint_512
|
||||
%In1 = OpTypeStruct %_arr_int_uint_512
|
||||
%_ptr_StorageBuffer_In1 = OpTypePointer StorageBuffer %In1
|
||||
%__1 = OpVariable %_ptr_StorageBuffer_In1 StorageBuffer
|
||||
%int_1 = OpConstant %int 1
|
||||
%_arr_int_uint_512_0 = OpTypeArray %int %uint_512
|
||||
%In0 = OpTypeStruct %_arr_int_uint_512_0
|
||||
%_ptr_StorageBuffer_In0 = OpTypePointer StorageBuffer %In0
|
||||
%__2 = OpVariable %_ptr_StorageBuffer_In0 StorageBuffer
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_1 %uint_1
|
||||
%main = OpFunction %void None %28
|
||||
%56 = OpLabel
|
||||
%base_index_in = OpVariable %_ptr_Function_uint Function
|
||||
%base_index_out = OpVariable %_ptr_Function_uint Function
|
||||
%index_in0 = OpVariable %_ptr_Function_int Function
|
||||
%index_in1 = OpVariable %_ptr_Function_int Function
|
||||
%index_out0 = OpVariable %_ptr_Function_int Function
|
||||
%index_out1 = OpVariable %_ptr_Function_int Function
|
||||
%condition_index = OpVariable %_ptr_Function_int Function
|
||||
%i = OpVariable %_ptr_Function_int Function
|
||||
%temp0 = OpVariable %_ptr_Function_int Function
|
||||
%temp1 = OpVariable %_ptr_Function_int Function
|
||||
%57 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0
|
||||
%58 = OpLoad %uint %57
|
||||
%59 = OpIMul %uint %uint_128 %58
|
||||
OpStore %base_index_in %59
|
||||
%60 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0
|
||||
%61 = OpLoad %uint %60
|
||||
%62 = OpIMul %uint %uint_256 %61
|
||||
OpStore %base_index_out %62
|
||||
OpStore %index_in0 %int_127
|
||||
OpStore %index_in1 %int_383
|
||||
OpStore %index_out0 %int_255
|
||||
OpStore %index_out1 %int_383
|
||||
OpStore %condition_index %int_0
|
||||
OpStore %i %int_0
|
||||
OpBranch %63
|
||||
%63 = OpLabel
|
||||
OpLoopMerge %64 %65 None
|
||||
OpBranch %66
|
||||
%66 = OpLabel
|
||||
%67 = OpLoad %int %i
|
||||
%68 = OpSLessThan %bool %67 %int_256
|
||||
OpBranchConditional %68 %69 %64
|
||||
%69 = OpLabel
|
||||
%70 = OpLoad %int %condition_index
|
||||
%71 = OpAccessChain %_ptr_StorageBuffer_int %_ %int_0 %70
|
||||
%72 = OpLoad %int %71
|
||||
%73 = OpIEqual %bool %72 %int_0
|
||||
OpSelectionMerge %74 None
|
||||
OpBranchConditional %73 %75 %76
|
||||
%75 = OpLabel
|
||||
%77 = OpLoad %uint %base_index_out
|
||||
%78 = OpLoad %int %index_out0
|
||||
%79 = OpBitcast %uint %78
|
||||
%80 = OpIAdd %uint %77 %79
|
||||
%81 = OpLoad %uint %base_index_in
|
||||
%82 = OpLoad %int %index_in0
|
||||
%83 = OpBitcast %uint %82
|
||||
%84 = OpIAdd %uint %81 %83
|
||||
%85 = OpAccessChain %_ptr_StorageBuffer_int %__1 %int_0 %84
|
||||
%86 = OpLoad %int %85
|
||||
%87 = OpAccessChain %_ptr_StorageBuffer_int %__0 %int_0 %80
|
||||
OpStore %87 %86
|
||||
%88 = OpLoad %int %index_out0
|
||||
%89 = OpISub %int %88 %int_1
|
||||
OpStore %index_out0 %89
|
||||
%90 = OpLoad %int %index_in1
|
||||
%91 = OpISub %int %90 %int_1
|
||||
OpStore %index_in1 %91
|
||||
OpBranch %74
|
||||
%76 = OpLabel
|
||||
%92 = OpLoad %uint %base_index_out
|
||||
%93 = OpLoad %int %index_out1
|
||||
%94 = OpBitcast %uint %93
|
||||
%95 = OpIAdd %uint %92 %94
|
||||
%96 = OpLoad %uint %base_index_in
|
||||
%97 = OpLoad %int %index_in1
|
||||
%98 = OpBitcast %uint %97
|
||||
%99 = OpIAdd %uint %96 %98
|
||||
%100 = OpAccessChain %_ptr_StorageBuffer_int %__2 %int_0 %99
|
||||
%101 = OpLoad %int %100
|
||||
%102 = OpAccessChain %_ptr_StorageBuffer_int %__0 %int_0 %95
|
||||
OpStore %102 %101
|
||||
%103 = OpLoad %int %index_out1
|
||||
%104 = OpISub %int %103 %int_1
|
||||
OpStore %index_out1 %104
|
||||
%105 = OpLoad %int %index_in1
|
||||
%106 = OpISub %int %105 %int_1
|
||||
OpStore %index_in1 %106
|
||||
OpBranch %74
|
||||
%74 = OpLabel
|
||||
%107 = OpLoad %int %condition_index
|
||||
%108 = OpIAdd %int %107 %int_1
|
||||
%109 = OpAccessChain %_ptr_StorageBuffer_int %_ %int_0 %108
|
||||
%110 = OpLoad %int %109
|
||||
%111 = OpLoad %int %condition_index
|
||||
%112 = OpIAdd %int %111 %110
|
||||
OpStore %condition_index %112
|
||||
%113 = OpLoad %int %index_in0
|
||||
OpStore %temp0 %113
|
||||
%114 = OpLoad %int %index_in1
|
||||
OpStore %index_in0 %114
|
||||
%115 = OpLoad %int %temp0
|
||||
OpStore %index_in1 %115
|
||||
%116 = OpLoad %int %index_out0
|
||||
OpStore %temp1 %116
|
||||
%117 = OpLoad %int %index_out1
|
||||
OpStore %index_out0 %117
|
||||
%118 = OpLoad %int %temp1
|
||||
OpStore %index_out1 %118
|
||||
OpBranch %65
|
||||
%65 = OpLabel
|
||||
%119 = OpLoad %int %i
|
||||
%120 = OpIAdd %int %119 %int_1
|
||||
OpStore %i %120
|
||||
OpBranch %63
|
||||
%64 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,69 @@
|
||||
static uint3 gl_WorkGroupID = uint3(0u, 0u, 0u);
|
||||
ByteAddressBuffer x_13 : register(t2, space0);
|
||||
RWByteAddressBuffer x_15 : register(u3, space0);
|
||||
ByteAddressBuffer x_17 : register(t1, space0);
|
||||
ByteAddressBuffer x_19 : register(t0, space0);
|
||||
|
||||
void main_1() {
|
||||
uint base_index_in = 0u;
|
||||
uint base_index_out = 0u;
|
||||
int index_in0 = 0;
|
||||
int index_in1 = 0;
|
||||
int index_out0 = 0;
|
||||
int index_out1 = 0;
|
||||
int condition_index = 0;
|
||||
int i = 0;
|
||||
int temp0 = 0;
|
||||
int temp1 = 0;
|
||||
const uint x_58 = gl_WorkGroupID.x;
|
||||
base_index_in = (128u * x_58);
|
||||
const uint x_61 = gl_WorkGroupID.x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
{
|
||||
for(; (i < 256); i = (i + 1)) {
|
||||
const int x_72 = asint(x_13.Load((4u * uint(condition_index))));
|
||||
if ((x_72 == 0)) {
|
||||
const uint x_77 = base_index_out;
|
||||
const int x_78 = index_out0;
|
||||
const int x_86 = asint(x_17.Load((4u * (base_index_in + asuint(index_in0)))));
|
||||
x_15.Store((4u * (x_77 + asuint(x_78))), asuint(x_86));
|
||||
index_out0 = (index_out0 - 1);
|
||||
index_in1 = (index_in1 - 1);
|
||||
} else {
|
||||
const uint x_92 = base_index_out;
|
||||
const int x_93 = index_out1;
|
||||
const int x_101 = asint(x_19.Load((4u * (base_index_in + asuint(index_in1)))));
|
||||
x_15.Store((4u * (x_92 + asuint(x_93))), asuint(x_101));
|
||||
index_out1 = (index_out1 - 1);
|
||||
index_in1 = (index_in1 - 1);
|
||||
}
|
||||
const int x_110 = asint(x_13.Load((4u * uint((condition_index + 1)))));
|
||||
condition_index = (condition_index + x_110);
|
||||
temp0 = index_in0;
|
||||
index_in0 = index_in1;
|
||||
index_in1 = temp0;
|
||||
temp1 = index_out0;
|
||||
index_out0 = index_out1;
|
||||
index_out1 = temp1;
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint3 gl_WorkGroupID_param : SV_GroupID;
|
||||
};
|
||||
|
||||
[numthreads(4, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
const uint3 gl_WorkGroupID_param = tint_symbol.gl_WorkGroupID_param;
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,108 @@
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ int arr[8];
|
||||
};
|
||||
struct In2 {
|
||||
/* 0x0000 */ tint_array_wrapper data_in2;
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
/* 0x0000 */ int arr[1024];
|
||||
};
|
||||
struct Out0 {
|
||||
/* 0x0000 */ tint_array_wrapper_1 data_out0;
|
||||
};
|
||||
struct tint_array_wrapper_2 {
|
||||
/* 0x0000 */ int arr[512];
|
||||
};
|
||||
struct In1 {
|
||||
/* 0x0000 */ tint_array_wrapper_2 data_in1;
|
||||
};
|
||||
struct In0 {
|
||||
/* 0x0000 */ tint_array_wrapper_2 data_in0;
|
||||
};
|
||||
|
||||
void main_1(const device In2& x_13, const device In1& x_17, device Out0& x_15, const device In0& x_19, thread uint3* const tint_symbol_2) {
|
||||
uint base_index_in = 0u;
|
||||
uint base_index_out = 0u;
|
||||
int index_in0 = 0;
|
||||
int index_in1 = 0;
|
||||
int index_out0 = 0;
|
||||
int index_out1 = 0;
|
||||
int condition_index = 0;
|
||||
int i = 0;
|
||||
int temp0 = 0;
|
||||
int temp1 = 0;
|
||||
uint const x_58 = (*(tint_symbol_2)).x;
|
||||
base_index_in = (128u * x_58);
|
||||
uint const x_61 = (*(tint_symbol_2)).x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
while (true) {
|
||||
int const x_67 = i;
|
||||
if ((x_67 < 256)) {
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
int const x_70 = condition_index;
|
||||
int const x_72 = x_13.data_in2.arr[x_70];
|
||||
if ((x_72 == 0)) {
|
||||
uint const x_77 = base_index_out;
|
||||
int const x_78 = index_out0;
|
||||
uint const x_81 = base_index_in;
|
||||
int const x_82 = index_in0;
|
||||
int const x_86 = x_17.data_in1.arr[(x_81 + as_type<uint>(x_82))];
|
||||
x_15.data_out0.arr[(x_77 + as_type<uint>(x_78))] = x_86;
|
||||
int const x_88 = index_out0;
|
||||
index_out0 = (x_88 - 1);
|
||||
int const x_90 = index_in1;
|
||||
index_in1 = (x_90 - 1);
|
||||
} else {
|
||||
uint const x_92 = base_index_out;
|
||||
int const x_93 = index_out1;
|
||||
uint const x_96 = base_index_in;
|
||||
int const x_97 = index_in1;
|
||||
int const x_101 = x_19.data_in0.arr[(x_96 + as_type<uint>(x_97))];
|
||||
x_15.data_out0.arr[(x_92 + as_type<uint>(x_93))] = x_101;
|
||||
int const x_103 = index_out1;
|
||||
index_out1 = (x_103 - 1);
|
||||
int const x_105 = index_in1;
|
||||
index_in1 = (x_105 - 1);
|
||||
}
|
||||
int const x_107 = condition_index;
|
||||
int const x_110 = x_13.data_in2.arr[(x_107 + 1)];
|
||||
int const x_111 = condition_index;
|
||||
condition_index = (x_111 + x_110);
|
||||
int const x_113 = index_in0;
|
||||
temp0 = x_113;
|
||||
int const x_114 = index_in1;
|
||||
index_in0 = x_114;
|
||||
int const x_115 = temp0;
|
||||
index_in1 = x_115;
|
||||
int const x_116 = index_out0;
|
||||
temp1 = x_116;
|
||||
int const x_117 = index_out1;
|
||||
index_out0 = x_117;
|
||||
int const x_118 = temp1;
|
||||
index_out1 = x_118;
|
||||
{
|
||||
int const x_119 = i;
|
||||
i = (x_119 + 1);
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint3 gl_WorkGroupID_param [[threadgroup_position_in_grid]], const device In2& x_13 [[buffer(2)]], const device In1& x_17 [[buffer(1)]], device Out0& x_15 [[buffer(3)]], const device In0& x_19 [[buffer(0)]]) {
|
||||
thread uint3 tint_symbol_3 = 0u;
|
||||
tint_symbol_3 = gl_WorkGroupID_param;
|
||||
main_1(x_13, x_17, x_15, x_19, &(tint_symbol_3));
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,224 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 127
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||
OpExecutionMode %main LocalSize 4 1 1
|
||||
OpName %gl_WorkGroupID "gl_WorkGroupID"
|
||||
OpName %In2 "In2"
|
||||
OpMemberName %In2 0 "data_in2"
|
||||
OpName %x_13 "x_13"
|
||||
OpName %Out0 "Out0"
|
||||
OpMemberName %Out0 0 "data_out0"
|
||||
OpName %x_15 "x_15"
|
||||
OpName %In1 "In1"
|
||||
OpMemberName %In1 0 "data_in1"
|
||||
OpName %x_17 "x_17"
|
||||
OpName %In0 "In0"
|
||||
OpMemberName %In0 0 "data_in0"
|
||||
OpName %x_19 "x_19"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %main_1 "main_1"
|
||||
OpName %base_index_in "base_index_in"
|
||||
OpName %base_index_out "base_index_out"
|
||||
OpName %index_in0 "index_in0"
|
||||
OpName %index_in1 "index_in1"
|
||||
OpName %index_out0 "index_out0"
|
||||
OpName %index_out1 "index_out1"
|
||||
OpName %condition_index "condition_index"
|
||||
OpName %i "i"
|
||||
OpName %temp0 "temp0"
|
||||
OpName %temp1 "temp1"
|
||||
OpName %main "main"
|
||||
OpDecorate %In2 Block
|
||||
OpMemberDecorate %In2 0 Offset 0
|
||||
OpDecorate %_arr_int_uint_8 ArrayStride 4
|
||||
OpDecorate %x_13 NonWritable
|
||||
OpDecorate %x_13 DescriptorSet 0
|
||||
OpDecorate %x_13 Binding 2
|
||||
OpDecorate %Out0 Block
|
||||
OpMemberDecorate %Out0 0 Offset 0
|
||||
OpDecorate %_arr_int_uint_1024 ArrayStride 4
|
||||
OpDecorate %x_15 DescriptorSet 0
|
||||
OpDecorate %x_15 Binding 3
|
||||
OpDecorate %In1 Block
|
||||
OpMemberDecorate %In1 0 Offset 0
|
||||
OpDecorate %_arr_int_uint_512 ArrayStride 4
|
||||
OpDecorate %x_17 NonWritable
|
||||
OpDecorate %x_17 DescriptorSet 0
|
||||
OpDecorate %x_17 Binding 1
|
||||
OpDecorate %In0 Block
|
||||
OpMemberDecorate %In0 0 Offset 0
|
||||
OpDecorate %x_19 NonWritable
|
||||
OpDecorate %x_19 DescriptorSet 0
|
||||
OpDecorate %x_19 Binding 0
|
||||
OpDecorate %tint_symbol BuiltIn WorkgroupId
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Private_v3uint = OpTypePointer Private %v3uint
|
||||
%5 = OpConstantNull %v3uint
|
||||
%gl_WorkGroupID = OpVariable %_ptr_Private_v3uint Private %5
|
||||
%int = OpTypeInt 32 1
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_int_uint_8 = OpTypeArray %int %uint_8
|
||||
%In2 = OpTypeStruct %_arr_int_uint_8
|
||||
%_ptr_StorageBuffer_In2 = OpTypePointer StorageBuffer %In2
|
||||
%x_13 = OpVariable %_ptr_StorageBuffer_In2 StorageBuffer
|
||||
%uint_1024 = OpConstant %uint 1024
|
||||
%_arr_int_uint_1024 = OpTypeArray %int %uint_1024
|
||||
%Out0 = OpTypeStruct %_arr_int_uint_1024
|
||||
%_ptr_StorageBuffer_Out0 = OpTypePointer StorageBuffer %Out0
|
||||
%x_15 = OpVariable %_ptr_StorageBuffer_Out0 StorageBuffer
|
||||
%uint_512 = OpConstant %uint 512
|
||||
%_arr_int_uint_512 = OpTypeArray %int %uint_512
|
||||
%In1 = OpTypeStruct %_arr_int_uint_512
|
||||
%_ptr_StorageBuffer_In1 = OpTypePointer StorageBuffer %In1
|
||||
%x_17 = OpVariable %_ptr_StorageBuffer_In1 StorageBuffer
|
||||
%In0 = OpTypeStruct %_arr_int_uint_512
|
||||
%_ptr_StorageBuffer_In0 = OpTypePointer StorageBuffer %In0
|
||||
%x_19 = OpVariable %_ptr_StorageBuffer_In0 StorageBuffer
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||
%void = OpTypeVoid
|
||||
%27 = OpTypeFunction %void
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%33 = OpConstantNull %uint
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%37 = OpConstantNull %int
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Private_uint = OpTypePointer Private %uint
|
||||
%uint_128 = OpConstant %uint 128
|
||||
%uint_256 = OpConstant %uint 256
|
||||
%int_127 = OpConstant %int 127
|
||||
%int_383 = OpConstant %int 383
|
||||
%int_255 = OpConstant %int 255
|
||||
%int_0 = OpConstant %int 0
|
||||
%int_256 = OpConstant %int 256
|
||||
%bool = OpTypeBool
|
||||
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
|
||||
%int_1 = OpConstant %int 1
|
||||
%main_1 = OpFunction %void None %27
|
||||
%30 = OpLabel
|
||||
%base_index_in = OpVariable %_ptr_Function_uint Function %33
|
||||
%base_index_out = OpVariable %_ptr_Function_uint Function %33
|
||||
%index_in0 = OpVariable %_ptr_Function_int Function %37
|
||||
%index_in1 = OpVariable %_ptr_Function_int Function %37
|
||||
%index_out0 = OpVariable %_ptr_Function_int Function %37
|
||||
%index_out1 = OpVariable %_ptr_Function_int Function %37
|
||||
%condition_index = OpVariable %_ptr_Function_int Function %37
|
||||
%i = OpVariable %_ptr_Function_int Function %37
|
||||
%temp0 = OpVariable %_ptr_Function_int Function %37
|
||||
%temp1 = OpVariable %_ptr_Function_int Function %37
|
||||
%47 = OpAccessChain %_ptr_Private_uint %gl_WorkGroupID %uint_0
|
||||
%48 = OpLoad %uint %47
|
||||
%50 = OpIMul %uint %uint_128 %48
|
||||
OpStore %base_index_in %50
|
||||
%51 = OpAccessChain %_ptr_Private_uint %gl_WorkGroupID %uint_0
|
||||
%52 = OpLoad %uint %51
|
||||
%54 = OpIMul %uint %uint_256 %52
|
||||
OpStore %base_index_out %54
|
||||
OpStore %index_in0 %int_127
|
||||
OpStore %index_in1 %int_383
|
||||
OpStore %index_out0 %int_255
|
||||
OpStore %index_out1 %int_383
|
||||
OpStore %condition_index %int_0
|
||||
OpStore %i %int_0
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
OpLoopMerge %60 %61 None
|
||||
OpBranch %62
|
||||
%62 = OpLabel
|
||||
%63 = OpLoad %int %i
|
||||
%65 = OpSLessThan %bool %63 %int_256
|
||||
OpSelectionMerge %67 None
|
||||
OpBranchConditional %65 %68 %69
|
||||
%68 = OpLabel
|
||||
OpBranch %67
|
||||
%69 = OpLabel
|
||||
OpBranch %60
|
||||
%67 = OpLabel
|
||||
%70 = OpLoad %int %condition_index
|
||||
%72 = OpAccessChain %_ptr_StorageBuffer_int %x_13 %uint_0 %70
|
||||
%73 = OpLoad %int %72
|
||||
%74 = OpIEqual %bool %73 %int_0
|
||||
OpSelectionMerge %75 None
|
||||
OpBranchConditional %74 %76 %77
|
||||
%76 = OpLabel
|
||||
%78 = OpLoad %uint %base_index_out
|
||||
%79 = OpLoad %int %index_out0
|
||||
%80 = OpLoad %uint %base_index_in
|
||||
%81 = OpLoad %int %index_in0
|
||||
%82 = OpBitcast %uint %81
|
||||
%83 = OpIAdd %uint %80 %82
|
||||
%84 = OpAccessChain %_ptr_StorageBuffer_int %x_17 %uint_0 %83
|
||||
%85 = OpLoad %int %84
|
||||
%86 = OpBitcast %uint %79
|
||||
%87 = OpIAdd %uint %78 %86
|
||||
%88 = OpAccessChain %_ptr_StorageBuffer_int %x_15 %uint_0 %87
|
||||
OpStore %88 %85
|
||||
%89 = OpLoad %int %index_out0
|
||||
%91 = OpISub %int %89 %int_1
|
||||
OpStore %index_out0 %91
|
||||
%92 = OpLoad %int %index_in1
|
||||
%93 = OpISub %int %92 %int_1
|
||||
OpStore %index_in1 %93
|
||||
OpBranch %75
|
||||
%77 = OpLabel
|
||||
%94 = OpLoad %uint %base_index_out
|
||||
%95 = OpLoad %int %index_out1
|
||||
%96 = OpLoad %uint %base_index_in
|
||||
%97 = OpLoad %int %index_in1
|
||||
%98 = OpBitcast %uint %97
|
||||
%99 = OpIAdd %uint %96 %98
|
||||
%100 = OpAccessChain %_ptr_StorageBuffer_int %x_19 %uint_0 %99
|
||||
%101 = OpLoad %int %100
|
||||
%102 = OpBitcast %uint %95
|
||||
%103 = OpIAdd %uint %94 %102
|
||||
%104 = OpAccessChain %_ptr_StorageBuffer_int %x_15 %uint_0 %103
|
||||
OpStore %104 %101
|
||||
%105 = OpLoad %int %index_out1
|
||||
%106 = OpISub %int %105 %int_1
|
||||
OpStore %index_out1 %106
|
||||
%107 = OpLoad %int %index_in1
|
||||
%108 = OpISub %int %107 %int_1
|
||||
OpStore %index_in1 %108
|
||||
OpBranch %75
|
||||
%75 = OpLabel
|
||||
%109 = OpLoad %int %condition_index
|
||||
%110 = OpIAdd %int %109 %int_1
|
||||
%111 = OpAccessChain %_ptr_StorageBuffer_int %x_13 %uint_0 %110
|
||||
%112 = OpLoad %int %111
|
||||
%113 = OpLoad %int %condition_index
|
||||
%114 = OpIAdd %int %113 %112
|
||||
OpStore %condition_index %114
|
||||
%115 = OpLoad %int %index_in0
|
||||
OpStore %temp0 %115
|
||||
%116 = OpLoad %int %index_in1
|
||||
OpStore %index_in0 %116
|
||||
%117 = OpLoad %int %temp0
|
||||
OpStore %index_in1 %117
|
||||
%118 = OpLoad %int %index_out0
|
||||
OpStore %temp1 %118
|
||||
%119 = OpLoad %int %index_out1
|
||||
OpStore %index_out0 %119
|
||||
%120 = OpLoad %int %temp1
|
||||
OpStore %index_out1 %120
|
||||
OpBranch %61
|
||||
%61 = OpLabel
|
||||
%121 = OpLoad %int %i
|
||||
%122 = OpIAdd %int %121 %int_1
|
||||
OpStore %i %122
|
||||
OpBranch %59
|
||||
%60 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %27
|
||||
%124 = OpLabel
|
||||
%125 = OpLoad %v3uint %tint_symbol
|
||||
OpStore %gl_WorkGroupID %125
|
||||
%126 = OpFunctionCall %void %main_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,120 @@
|
||||
type Arr = [[stride(4)]] array<i32, 8>;
|
||||
|
||||
[[block]]
|
||||
struct In2 {
|
||||
data_in2 : Arr;
|
||||
};
|
||||
|
||||
type Arr_1 = [[stride(4)]] array<i32, 1024>;
|
||||
|
||||
[[block]]
|
||||
struct Out0 {
|
||||
data_out0 : Arr_1;
|
||||
};
|
||||
|
||||
type Arr_2 = [[stride(4)]] array<i32, 512>;
|
||||
|
||||
type Arr_3 = [[stride(4)]] array<i32, 512>;
|
||||
|
||||
[[block]]
|
||||
struct In1 {
|
||||
data_in1 : Arr_3;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct In0 {
|
||||
data_in0 : Arr_3;
|
||||
};
|
||||
|
||||
var<private> gl_WorkGroupID : vec3<u32>;
|
||||
|
||||
[[group(0), binding(2)]] var<storage, read> x_13 : In2;
|
||||
|
||||
[[group(0), binding(3)]] var<storage, read_write> x_15 : Out0;
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> x_17 : In1;
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read> x_19 : In0;
|
||||
|
||||
fn main_1() {
|
||||
var base_index_in : u32;
|
||||
var base_index_out : u32;
|
||||
var index_in0 : i32;
|
||||
var index_in1 : i32;
|
||||
var index_out0 : i32;
|
||||
var index_out1 : i32;
|
||||
var condition_index : i32;
|
||||
var i : i32;
|
||||
var temp0 : i32;
|
||||
var temp1 : i32;
|
||||
let x_58 : u32 = gl_WorkGroupID.x;
|
||||
base_index_in = (128u * x_58);
|
||||
let x_61 : u32 = gl_WorkGroupID.x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
loop {
|
||||
let x_67 : i32 = i;
|
||||
if ((x_67 < 256)) {
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
let x_70 : i32 = condition_index;
|
||||
let x_72 : i32 = x_13.data_in2[x_70];
|
||||
if ((x_72 == 0)) {
|
||||
let x_77 : u32 = base_index_out;
|
||||
let x_78 : i32 = index_out0;
|
||||
let x_81 : u32 = base_index_in;
|
||||
let x_82 : i32 = index_in0;
|
||||
let x_86 : i32 = x_17.data_in1[(x_81 + bitcast<u32>(x_82))];
|
||||
x_15.data_out0[(x_77 + bitcast<u32>(x_78))] = x_86;
|
||||
let x_88 : i32 = index_out0;
|
||||
index_out0 = (x_88 - 1);
|
||||
let x_90 : i32 = index_in1;
|
||||
index_in1 = (x_90 - 1);
|
||||
} else {
|
||||
let x_92 : u32 = base_index_out;
|
||||
let x_93 : i32 = index_out1;
|
||||
let x_96 : u32 = base_index_in;
|
||||
let x_97 : i32 = index_in1;
|
||||
let x_101 : i32 = x_19.data_in0[(x_96 + bitcast<u32>(x_97))];
|
||||
x_15.data_out0[(x_92 + bitcast<u32>(x_93))] = x_101;
|
||||
let x_103 : i32 = index_out1;
|
||||
index_out1 = (x_103 - 1);
|
||||
let x_105 : i32 = index_in1;
|
||||
index_in1 = (x_105 - 1);
|
||||
}
|
||||
let x_107 : i32 = condition_index;
|
||||
let x_110 : i32 = x_13.data_in2[(x_107 + 1)];
|
||||
let x_111 : i32 = condition_index;
|
||||
condition_index = (x_111 + x_110);
|
||||
let x_113 : i32 = index_in0;
|
||||
temp0 = x_113;
|
||||
let x_114 : i32 = index_in1;
|
||||
index_in0 = x_114;
|
||||
let x_115 : i32 = temp0;
|
||||
index_in1 = x_115;
|
||||
let x_116 : i32 = index_out0;
|
||||
temp1 = x_116;
|
||||
let x_117 : i32 = index_out1;
|
||||
index_out0 = x_117;
|
||||
let x_118 : i32 = temp1;
|
||||
index_out1 = x_118;
|
||||
|
||||
continuing {
|
||||
let x_119 : i32 = i;
|
||||
i = (x_119 + 1);
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(4, 1, 1)]]
|
||||
fn main([[builtin(workgroup_id)]] gl_WorkGroupID_param : vec3<u32>) {
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
}
|
||||
@@ -0,0 +1,120 @@
|
||||
type Arr = [[stride(4)]] array<i32, 8>;
|
||||
|
||||
[[block]]
|
||||
struct In2 {
|
||||
data_in2 : Arr;
|
||||
};
|
||||
|
||||
type Arr_1 = [[stride(4)]] array<i32, 1024>;
|
||||
|
||||
[[block]]
|
||||
struct Out0 {
|
||||
data_out0 : Arr_1;
|
||||
};
|
||||
|
||||
type Arr_2 = [[stride(4)]] array<i32, 512>;
|
||||
|
||||
type Arr_3 = [[stride(4)]] array<i32, 512>;
|
||||
|
||||
[[block]]
|
||||
struct In1 {
|
||||
data_in1 : Arr_3;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct In0 {
|
||||
data_in0 : Arr_3;
|
||||
};
|
||||
|
||||
var<private> gl_WorkGroupID : vec3<u32>;
|
||||
|
||||
[[group(0), binding(2)]] var<storage, read> x_13 : In2;
|
||||
|
||||
[[group(0), binding(3)]] var<storage, read_write> x_15 : Out0;
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> x_17 : In1;
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read> x_19 : In0;
|
||||
|
||||
fn main_1() {
|
||||
var base_index_in : u32;
|
||||
var base_index_out : u32;
|
||||
var index_in0 : i32;
|
||||
var index_in1 : i32;
|
||||
var index_out0 : i32;
|
||||
var index_out1 : i32;
|
||||
var condition_index : i32;
|
||||
var i : i32;
|
||||
var temp0 : i32;
|
||||
var temp1 : i32;
|
||||
let x_58 : u32 = gl_WorkGroupID.x;
|
||||
base_index_in = (128u * x_58);
|
||||
let x_61 : u32 = gl_WorkGroupID.x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
loop {
|
||||
let x_67 : i32 = i;
|
||||
if ((x_67 < 256)) {
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
let x_70 : i32 = condition_index;
|
||||
let x_72 : i32 = x_13.data_in2[x_70];
|
||||
if ((x_72 == 0)) {
|
||||
let x_77 : u32 = base_index_out;
|
||||
let x_78 : i32 = index_out0;
|
||||
let x_81 : u32 = base_index_in;
|
||||
let x_82 : i32 = index_in0;
|
||||
let x_86 : i32 = x_17.data_in1[(x_81 + bitcast<u32>(x_82))];
|
||||
x_15.data_out0[(x_77 + bitcast<u32>(x_78))] = x_86;
|
||||
let x_88 : i32 = index_out0;
|
||||
index_out0 = (x_88 - 1);
|
||||
let x_90 : i32 = index_in1;
|
||||
index_in1 = (x_90 - 1);
|
||||
} else {
|
||||
let x_92 : u32 = base_index_out;
|
||||
let x_93 : i32 = index_out1;
|
||||
let x_96 : u32 = base_index_in;
|
||||
let x_97 : i32 = index_in1;
|
||||
let x_101 : i32 = x_19.data_in0[(x_96 + bitcast<u32>(x_97))];
|
||||
x_15.data_out0[(x_92 + bitcast<u32>(x_93))] = x_101;
|
||||
let x_103 : i32 = index_out1;
|
||||
index_out1 = (x_103 - 1);
|
||||
let x_105 : i32 = index_in1;
|
||||
index_in1 = (x_105 - 1);
|
||||
}
|
||||
let x_107 : i32 = condition_index;
|
||||
let x_110 : i32 = x_13.data_in2[(x_107 + 1)];
|
||||
let x_111 : i32 = condition_index;
|
||||
condition_index = (x_111 + x_110);
|
||||
let x_113 : i32 = index_in0;
|
||||
temp0 = x_113;
|
||||
let x_114 : i32 = index_in1;
|
||||
index_in0 = x_114;
|
||||
let x_115 : i32 = temp0;
|
||||
index_in1 = x_115;
|
||||
let x_116 : i32 = index_out0;
|
||||
temp1 = x_116;
|
||||
let x_117 : i32 = index_out1;
|
||||
index_out0 = x_117;
|
||||
let x_118 : i32 = temp1;
|
||||
index_out1 = x_118;
|
||||
|
||||
continuing {
|
||||
let x_119 : i32 = i;
|
||||
i = (x_119 + 1);
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(4, 1, 1)]]
|
||||
fn main([[builtin(workgroup_id)]] gl_WorkGroupID_param : vec3<u32>) {
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
}
|
||||
@@ -0,0 +1,69 @@
|
||||
static uint3 gl_WorkGroupID = uint3(0u, 0u, 0u);
|
||||
ByteAddressBuffer x_13 : register(t2, space0);
|
||||
RWByteAddressBuffer x_15 : register(u3, space0);
|
||||
ByteAddressBuffer x_17 : register(t1, space0);
|
||||
ByteAddressBuffer x_19 : register(t0, space0);
|
||||
|
||||
void main_1() {
|
||||
uint base_index_in = 0u;
|
||||
uint base_index_out = 0u;
|
||||
int index_in0 = 0;
|
||||
int index_in1 = 0;
|
||||
int index_out0 = 0;
|
||||
int index_out1 = 0;
|
||||
int condition_index = 0;
|
||||
int i = 0;
|
||||
int temp0 = 0;
|
||||
int temp1 = 0;
|
||||
const uint x_58 = gl_WorkGroupID.x;
|
||||
base_index_in = (128u * x_58);
|
||||
const uint x_61 = gl_WorkGroupID.x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
{
|
||||
for(; (i < 256); i = (i + 1)) {
|
||||
const int x_72 = asint(x_13.Load((4u * uint(condition_index))));
|
||||
if ((x_72 == 0)) {
|
||||
const uint x_77 = base_index_out;
|
||||
const int x_78 = index_out0;
|
||||
const int x_86 = asint(x_17.Load((4u * (base_index_in + asuint(index_in0)))));
|
||||
x_15.Store((4u * (x_77 + asuint(x_78))), asuint(x_86));
|
||||
index_out0 = (index_out0 - 1);
|
||||
index_in1 = (index_in1 - 1);
|
||||
} else {
|
||||
const uint x_92 = base_index_out;
|
||||
const int x_93 = index_out1;
|
||||
const int x_101 = asint(x_19.Load((4u * (base_index_in + asuint(index_in1)))));
|
||||
x_15.Store((4u * (x_92 + asuint(x_93))), asuint(x_101));
|
||||
index_out1 = (index_out1 - 1);
|
||||
index_in1 = (index_in1 - 1);
|
||||
}
|
||||
const int x_110 = asint(x_13.Load((4u * uint((condition_index + 1)))));
|
||||
condition_index = (condition_index + x_110);
|
||||
temp0 = index_in0;
|
||||
index_in0 = index_in1;
|
||||
index_in1 = temp0;
|
||||
temp1 = index_out0;
|
||||
index_out0 = index_out1;
|
||||
index_out1 = temp1;
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint3 gl_WorkGroupID_param : SV_GroupID;
|
||||
};
|
||||
|
||||
[numthreads(4, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
const uint3 gl_WorkGroupID_param = tint_symbol.gl_WorkGroupID_param;
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
return;
|
||||
}
|
||||
@@ -0,0 +1,108 @@
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ int arr[8];
|
||||
};
|
||||
struct In2 {
|
||||
/* 0x0000 */ tint_array_wrapper data_in2;
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
/* 0x0000 */ int arr[1024];
|
||||
};
|
||||
struct Out0 {
|
||||
/* 0x0000 */ tint_array_wrapper_1 data_out0;
|
||||
};
|
||||
struct tint_array_wrapper_2 {
|
||||
/* 0x0000 */ int arr[512];
|
||||
};
|
||||
struct In1 {
|
||||
/* 0x0000 */ tint_array_wrapper_2 data_in1;
|
||||
};
|
||||
struct In0 {
|
||||
/* 0x0000 */ tint_array_wrapper_2 data_in0;
|
||||
};
|
||||
|
||||
void main_1(const device In2& x_13, const device In1& x_17, device Out0& x_15, const device In0& x_19, thread uint3* const tint_symbol_2) {
|
||||
uint base_index_in = 0u;
|
||||
uint base_index_out = 0u;
|
||||
int index_in0 = 0;
|
||||
int index_in1 = 0;
|
||||
int index_out0 = 0;
|
||||
int index_out1 = 0;
|
||||
int condition_index = 0;
|
||||
int i = 0;
|
||||
int temp0 = 0;
|
||||
int temp1 = 0;
|
||||
uint const x_58 = (*(tint_symbol_2)).x;
|
||||
base_index_in = (128u * x_58);
|
||||
uint const x_61 = (*(tint_symbol_2)).x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
while (true) {
|
||||
int const x_67 = i;
|
||||
if ((x_67 < 256)) {
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
int const x_70 = condition_index;
|
||||
int const x_72 = x_13.data_in2.arr[x_70];
|
||||
if ((x_72 == 0)) {
|
||||
uint const x_77 = base_index_out;
|
||||
int const x_78 = index_out0;
|
||||
uint const x_81 = base_index_in;
|
||||
int const x_82 = index_in0;
|
||||
int const x_86 = x_17.data_in1.arr[(x_81 + as_type<uint>(x_82))];
|
||||
x_15.data_out0.arr[(x_77 + as_type<uint>(x_78))] = x_86;
|
||||
int const x_88 = index_out0;
|
||||
index_out0 = (x_88 - 1);
|
||||
int const x_90 = index_in1;
|
||||
index_in1 = (x_90 - 1);
|
||||
} else {
|
||||
uint const x_92 = base_index_out;
|
||||
int const x_93 = index_out1;
|
||||
uint const x_96 = base_index_in;
|
||||
int const x_97 = index_in1;
|
||||
int const x_101 = x_19.data_in0.arr[(x_96 + as_type<uint>(x_97))];
|
||||
x_15.data_out0.arr[(x_92 + as_type<uint>(x_93))] = x_101;
|
||||
int const x_103 = index_out1;
|
||||
index_out1 = (x_103 - 1);
|
||||
int const x_105 = index_in1;
|
||||
index_in1 = (x_105 - 1);
|
||||
}
|
||||
int const x_107 = condition_index;
|
||||
int const x_110 = x_13.data_in2.arr[(x_107 + 1)];
|
||||
int const x_111 = condition_index;
|
||||
condition_index = (x_111 + x_110);
|
||||
int const x_113 = index_in0;
|
||||
temp0 = x_113;
|
||||
int const x_114 = index_in1;
|
||||
index_in0 = x_114;
|
||||
int const x_115 = temp0;
|
||||
index_in1 = x_115;
|
||||
int const x_116 = index_out0;
|
||||
temp1 = x_116;
|
||||
int const x_117 = index_out1;
|
||||
index_out0 = x_117;
|
||||
int const x_118 = temp1;
|
||||
index_out1 = x_118;
|
||||
{
|
||||
int const x_119 = i;
|
||||
i = (x_119 + 1);
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint3 gl_WorkGroupID_param [[threadgroup_position_in_grid]], const device In2& x_13 [[buffer(2)]], const device In1& x_17 [[buffer(1)]], device Out0& x_15 [[buffer(3)]], const device In0& x_19 [[buffer(0)]]) {
|
||||
thread uint3 tint_symbol_3 = 0u;
|
||||
tint_symbol_3 = gl_WorkGroupID_param;
|
||||
main_1(x_13, x_17, x_15, x_19, &(tint_symbol_3));
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,224 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 127
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||
OpExecutionMode %main LocalSize 4 1 1
|
||||
OpName %gl_WorkGroupID "gl_WorkGroupID"
|
||||
OpName %In2 "In2"
|
||||
OpMemberName %In2 0 "data_in2"
|
||||
OpName %x_13 "x_13"
|
||||
OpName %Out0 "Out0"
|
||||
OpMemberName %Out0 0 "data_out0"
|
||||
OpName %x_15 "x_15"
|
||||
OpName %In1 "In1"
|
||||
OpMemberName %In1 0 "data_in1"
|
||||
OpName %x_17 "x_17"
|
||||
OpName %In0 "In0"
|
||||
OpMemberName %In0 0 "data_in0"
|
||||
OpName %x_19 "x_19"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %main_1 "main_1"
|
||||
OpName %base_index_in "base_index_in"
|
||||
OpName %base_index_out "base_index_out"
|
||||
OpName %index_in0 "index_in0"
|
||||
OpName %index_in1 "index_in1"
|
||||
OpName %index_out0 "index_out0"
|
||||
OpName %index_out1 "index_out1"
|
||||
OpName %condition_index "condition_index"
|
||||
OpName %i "i"
|
||||
OpName %temp0 "temp0"
|
||||
OpName %temp1 "temp1"
|
||||
OpName %main "main"
|
||||
OpDecorate %In2 Block
|
||||
OpMemberDecorate %In2 0 Offset 0
|
||||
OpDecorate %_arr_int_uint_8 ArrayStride 4
|
||||
OpDecorate %x_13 NonWritable
|
||||
OpDecorate %x_13 DescriptorSet 0
|
||||
OpDecorate %x_13 Binding 2
|
||||
OpDecorate %Out0 Block
|
||||
OpMemberDecorate %Out0 0 Offset 0
|
||||
OpDecorate %_arr_int_uint_1024 ArrayStride 4
|
||||
OpDecorate %x_15 DescriptorSet 0
|
||||
OpDecorate %x_15 Binding 3
|
||||
OpDecorate %In1 Block
|
||||
OpMemberDecorate %In1 0 Offset 0
|
||||
OpDecorate %_arr_int_uint_512 ArrayStride 4
|
||||
OpDecorate %x_17 NonWritable
|
||||
OpDecorate %x_17 DescriptorSet 0
|
||||
OpDecorate %x_17 Binding 1
|
||||
OpDecorate %In0 Block
|
||||
OpMemberDecorate %In0 0 Offset 0
|
||||
OpDecorate %x_19 NonWritable
|
||||
OpDecorate %x_19 DescriptorSet 0
|
||||
OpDecorate %x_19 Binding 0
|
||||
OpDecorate %tint_symbol BuiltIn WorkgroupId
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Private_v3uint = OpTypePointer Private %v3uint
|
||||
%5 = OpConstantNull %v3uint
|
||||
%gl_WorkGroupID = OpVariable %_ptr_Private_v3uint Private %5
|
||||
%int = OpTypeInt 32 1
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_int_uint_8 = OpTypeArray %int %uint_8
|
||||
%In2 = OpTypeStruct %_arr_int_uint_8
|
||||
%_ptr_StorageBuffer_In2 = OpTypePointer StorageBuffer %In2
|
||||
%x_13 = OpVariable %_ptr_StorageBuffer_In2 StorageBuffer
|
||||
%uint_1024 = OpConstant %uint 1024
|
||||
%_arr_int_uint_1024 = OpTypeArray %int %uint_1024
|
||||
%Out0 = OpTypeStruct %_arr_int_uint_1024
|
||||
%_ptr_StorageBuffer_Out0 = OpTypePointer StorageBuffer %Out0
|
||||
%x_15 = OpVariable %_ptr_StorageBuffer_Out0 StorageBuffer
|
||||
%uint_512 = OpConstant %uint 512
|
||||
%_arr_int_uint_512 = OpTypeArray %int %uint_512
|
||||
%In1 = OpTypeStruct %_arr_int_uint_512
|
||||
%_ptr_StorageBuffer_In1 = OpTypePointer StorageBuffer %In1
|
||||
%x_17 = OpVariable %_ptr_StorageBuffer_In1 StorageBuffer
|
||||
%In0 = OpTypeStruct %_arr_int_uint_512
|
||||
%_ptr_StorageBuffer_In0 = OpTypePointer StorageBuffer %In0
|
||||
%x_19 = OpVariable %_ptr_StorageBuffer_In0 StorageBuffer
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||
%void = OpTypeVoid
|
||||
%27 = OpTypeFunction %void
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%33 = OpConstantNull %uint
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%37 = OpConstantNull %int
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Private_uint = OpTypePointer Private %uint
|
||||
%uint_128 = OpConstant %uint 128
|
||||
%uint_256 = OpConstant %uint 256
|
||||
%int_127 = OpConstant %int 127
|
||||
%int_383 = OpConstant %int 383
|
||||
%int_255 = OpConstant %int 255
|
||||
%int_0 = OpConstant %int 0
|
||||
%int_256 = OpConstant %int 256
|
||||
%bool = OpTypeBool
|
||||
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
|
||||
%int_1 = OpConstant %int 1
|
||||
%main_1 = OpFunction %void None %27
|
||||
%30 = OpLabel
|
||||
%base_index_in = OpVariable %_ptr_Function_uint Function %33
|
||||
%base_index_out = OpVariable %_ptr_Function_uint Function %33
|
||||
%index_in0 = OpVariable %_ptr_Function_int Function %37
|
||||
%index_in1 = OpVariable %_ptr_Function_int Function %37
|
||||
%index_out0 = OpVariable %_ptr_Function_int Function %37
|
||||
%index_out1 = OpVariable %_ptr_Function_int Function %37
|
||||
%condition_index = OpVariable %_ptr_Function_int Function %37
|
||||
%i = OpVariable %_ptr_Function_int Function %37
|
||||
%temp0 = OpVariable %_ptr_Function_int Function %37
|
||||
%temp1 = OpVariable %_ptr_Function_int Function %37
|
||||
%47 = OpAccessChain %_ptr_Private_uint %gl_WorkGroupID %uint_0
|
||||
%48 = OpLoad %uint %47
|
||||
%50 = OpIMul %uint %uint_128 %48
|
||||
OpStore %base_index_in %50
|
||||
%51 = OpAccessChain %_ptr_Private_uint %gl_WorkGroupID %uint_0
|
||||
%52 = OpLoad %uint %51
|
||||
%54 = OpIMul %uint %uint_256 %52
|
||||
OpStore %base_index_out %54
|
||||
OpStore %index_in0 %int_127
|
||||
OpStore %index_in1 %int_383
|
||||
OpStore %index_out0 %int_255
|
||||
OpStore %index_out1 %int_383
|
||||
OpStore %condition_index %int_0
|
||||
OpStore %i %int_0
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
OpLoopMerge %60 %61 None
|
||||
OpBranch %62
|
||||
%62 = OpLabel
|
||||
%63 = OpLoad %int %i
|
||||
%65 = OpSLessThan %bool %63 %int_256
|
||||
OpSelectionMerge %67 None
|
||||
OpBranchConditional %65 %68 %69
|
||||
%68 = OpLabel
|
||||
OpBranch %67
|
||||
%69 = OpLabel
|
||||
OpBranch %60
|
||||
%67 = OpLabel
|
||||
%70 = OpLoad %int %condition_index
|
||||
%72 = OpAccessChain %_ptr_StorageBuffer_int %x_13 %uint_0 %70
|
||||
%73 = OpLoad %int %72
|
||||
%74 = OpIEqual %bool %73 %int_0
|
||||
OpSelectionMerge %75 None
|
||||
OpBranchConditional %74 %76 %77
|
||||
%76 = OpLabel
|
||||
%78 = OpLoad %uint %base_index_out
|
||||
%79 = OpLoad %int %index_out0
|
||||
%80 = OpLoad %uint %base_index_in
|
||||
%81 = OpLoad %int %index_in0
|
||||
%82 = OpBitcast %uint %81
|
||||
%83 = OpIAdd %uint %80 %82
|
||||
%84 = OpAccessChain %_ptr_StorageBuffer_int %x_17 %uint_0 %83
|
||||
%85 = OpLoad %int %84
|
||||
%86 = OpBitcast %uint %79
|
||||
%87 = OpIAdd %uint %78 %86
|
||||
%88 = OpAccessChain %_ptr_StorageBuffer_int %x_15 %uint_0 %87
|
||||
OpStore %88 %85
|
||||
%89 = OpLoad %int %index_out0
|
||||
%91 = OpISub %int %89 %int_1
|
||||
OpStore %index_out0 %91
|
||||
%92 = OpLoad %int %index_in1
|
||||
%93 = OpISub %int %92 %int_1
|
||||
OpStore %index_in1 %93
|
||||
OpBranch %75
|
||||
%77 = OpLabel
|
||||
%94 = OpLoad %uint %base_index_out
|
||||
%95 = OpLoad %int %index_out1
|
||||
%96 = OpLoad %uint %base_index_in
|
||||
%97 = OpLoad %int %index_in1
|
||||
%98 = OpBitcast %uint %97
|
||||
%99 = OpIAdd %uint %96 %98
|
||||
%100 = OpAccessChain %_ptr_StorageBuffer_int %x_19 %uint_0 %99
|
||||
%101 = OpLoad %int %100
|
||||
%102 = OpBitcast %uint %95
|
||||
%103 = OpIAdd %uint %94 %102
|
||||
%104 = OpAccessChain %_ptr_StorageBuffer_int %x_15 %uint_0 %103
|
||||
OpStore %104 %101
|
||||
%105 = OpLoad %int %index_out1
|
||||
%106 = OpISub %int %105 %int_1
|
||||
OpStore %index_out1 %106
|
||||
%107 = OpLoad %int %index_in1
|
||||
%108 = OpISub %int %107 %int_1
|
||||
OpStore %index_in1 %108
|
||||
OpBranch %75
|
||||
%75 = OpLabel
|
||||
%109 = OpLoad %int %condition_index
|
||||
%110 = OpIAdd %int %109 %int_1
|
||||
%111 = OpAccessChain %_ptr_StorageBuffer_int %x_13 %uint_0 %110
|
||||
%112 = OpLoad %int %111
|
||||
%113 = OpLoad %int %condition_index
|
||||
%114 = OpIAdd %int %113 %112
|
||||
OpStore %condition_index %114
|
||||
%115 = OpLoad %int %index_in0
|
||||
OpStore %temp0 %115
|
||||
%116 = OpLoad %int %index_in1
|
||||
OpStore %index_in0 %116
|
||||
%117 = OpLoad %int %temp0
|
||||
OpStore %index_in1 %117
|
||||
%118 = OpLoad %int %index_out0
|
||||
OpStore %temp1 %118
|
||||
%119 = OpLoad %int %index_out1
|
||||
OpStore %index_out0 %119
|
||||
%120 = OpLoad %int %temp1
|
||||
OpStore %index_out1 %120
|
||||
OpBranch %61
|
||||
%61 = OpLabel
|
||||
%121 = OpLoad %int %i
|
||||
%122 = OpIAdd %int %121 %int_1
|
||||
OpStore %i %122
|
||||
OpBranch %59
|
||||
%60 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %27
|
||||
%124 = OpLabel
|
||||
%125 = OpLoad %v3uint %tint_symbol
|
||||
OpStore %gl_WorkGroupID %125
|
||||
%126 = OpFunctionCall %void %main_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -0,0 +1,120 @@
|
||||
type Arr = [[stride(4)]] array<i32, 8>;
|
||||
|
||||
[[block]]
|
||||
struct In2 {
|
||||
data_in2 : Arr;
|
||||
};
|
||||
|
||||
type Arr_1 = [[stride(4)]] array<i32, 1024>;
|
||||
|
||||
[[block]]
|
||||
struct Out0 {
|
||||
data_out0 : Arr_1;
|
||||
};
|
||||
|
||||
type Arr_2 = [[stride(4)]] array<i32, 512>;
|
||||
|
||||
type Arr_3 = [[stride(4)]] array<i32, 512>;
|
||||
|
||||
[[block]]
|
||||
struct In1 {
|
||||
data_in1 : Arr_3;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct In0 {
|
||||
data_in0 : Arr_3;
|
||||
};
|
||||
|
||||
var<private> gl_WorkGroupID : vec3<u32>;
|
||||
|
||||
[[group(0), binding(2)]] var<storage, read> x_13 : In2;
|
||||
|
||||
[[group(0), binding(3)]] var<storage, read_write> x_15 : Out0;
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> x_17 : In1;
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read> x_19 : In0;
|
||||
|
||||
fn main_1() {
|
||||
var base_index_in : u32;
|
||||
var base_index_out : u32;
|
||||
var index_in0 : i32;
|
||||
var index_in1 : i32;
|
||||
var index_out0 : i32;
|
||||
var index_out1 : i32;
|
||||
var condition_index : i32;
|
||||
var i : i32;
|
||||
var temp0 : i32;
|
||||
var temp1 : i32;
|
||||
let x_58 : u32 = gl_WorkGroupID.x;
|
||||
base_index_in = (128u * x_58);
|
||||
let x_61 : u32 = gl_WorkGroupID.x;
|
||||
base_index_out = (256u * x_61);
|
||||
index_in0 = 127;
|
||||
index_in1 = 383;
|
||||
index_out0 = 255;
|
||||
index_out1 = 383;
|
||||
condition_index = 0;
|
||||
i = 0;
|
||||
loop {
|
||||
let x_67 : i32 = i;
|
||||
if ((x_67 < 256)) {
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
let x_70 : i32 = condition_index;
|
||||
let x_72 : i32 = x_13.data_in2[x_70];
|
||||
if ((x_72 == 0)) {
|
||||
let x_77 : u32 = base_index_out;
|
||||
let x_78 : i32 = index_out0;
|
||||
let x_81 : u32 = base_index_in;
|
||||
let x_82 : i32 = index_in0;
|
||||
let x_86 : i32 = x_17.data_in1[(x_81 + bitcast<u32>(x_82))];
|
||||
x_15.data_out0[(x_77 + bitcast<u32>(x_78))] = x_86;
|
||||
let x_88 : i32 = index_out0;
|
||||
index_out0 = (x_88 - 1);
|
||||
let x_90 : i32 = index_in1;
|
||||
index_in1 = (x_90 - 1);
|
||||
} else {
|
||||
let x_92 : u32 = base_index_out;
|
||||
let x_93 : i32 = index_out1;
|
||||
let x_96 : u32 = base_index_in;
|
||||
let x_97 : i32 = index_in1;
|
||||
let x_101 : i32 = x_19.data_in0[(x_96 + bitcast<u32>(x_97))];
|
||||
x_15.data_out0[(x_92 + bitcast<u32>(x_93))] = x_101;
|
||||
let x_103 : i32 = index_out1;
|
||||
index_out1 = (x_103 - 1);
|
||||
let x_105 : i32 = index_in1;
|
||||
index_in1 = (x_105 - 1);
|
||||
}
|
||||
let x_107 : i32 = condition_index;
|
||||
let x_110 : i32 = x_13.data_in2[(x_107 + 1)];
|
||||
let x_111 : i32 = condition_index;
|
||||
condition_index = (x_111 + x_110);
|
||||
let x_113 : i32 = index_in0;
|
||||
temp0 = x_113;
|
||||
let x_114 : i32 = index_in1;
|
||||
index_in0 = x_114;
|
||||
let x_115 : i32 = temp0;
|
||||
index_in1 = x_115;
|
||||
let x_116 : i32 = index_out0;
|
||||
temp1 = x_116;
|
||||
let x_117 : i32 = index_out1;
|
||||
index_out0 = x_117;
|
||||
let x_118 : i32 = temp1;
|
||||
index_out1 = x_118;
|
||||
|
||||
continuing {
|
||||
let x_119 : i32 = i;
|
||||
i = (x_119 + 1);
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(4, 1, 1)]]
|
||||
fn main([[builtin(workgroup_id)]] gl_WorkGroupID_param : vec3<u32>) {
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
}
|
||||
Reference in New Issue
Block a user