mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-12-10 22:17:51 +00:00
test: Remove many expected files
For these tests, we only really care that we can successfully consume them and generate valid output for each backend. Having the expected files in the tree generates significant churn for any change to how we generate backend code, which makes it hard to inspect diffs. Change-Id: Ic98c248081144c0fb1791f1303eaf6d459548e3d Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62720 Reviewed-by: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: James Price <jrprice@google.com>
This commit is contained in:
James Price
committed by
Tint LUCI CQ
Tint LUCI CQ
parent
568136dd10
commit
87cce20f67
@@ -1,72 +0,0 @@
|
||||
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;
|
||||
};
|
||||
|
||||
void main_inner(uint3 gl_WorkGroupID_param) {
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
}
|
||||
|
||||
[numthreads(4, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
main_inner(tint_symbol.gl_WorkGroupID_param);
|
||||
return;
|
||||
}
|
||||
@@ -1,112 +0,0 @@
|
||||
#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_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;
|
||||
uint const x_58 = (*(tint_symbol_1)).x;
|
||||
base_index_in = (128u * x_58);
|
||||
uint const x_61 = (*(tint_symbol_1)).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 = as_type<int>((as_type<uint>(x_88) - as_type<uint>(1)));
|
||||
int const x_90 = index_in1;
|
||||
index_in1 = as_type<int>((as_type<uint>(x_90) - as_type<uint>(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 = as_type<int>((as_type<uint>(x_103) - as_type<uint>(1)));
|
||||
int const x_105 = index_in1;
|
||||
index_in1 = as_type<int>((as_type<uint>(x_105) - as_type<uint>(1)));
|
||||
}
|
||||
int const x_107 = condition_index;
|
||||
int const x_110 = x_13.data_in2.arr[as_type<int>((as_type<uint>(x_107) + as_type<uint>(1)))];
|
||||
int const x_111 = condition_index;
|
||||
condition_index = as_type<int>((as_type<uint>(x_111) + as_type<uint>(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 = as_type<int>((as_type<uint>(x_119) + as_type<uint>(1)));
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
void tint_symbol_inner(const device In2& x_13, const device In1& x_17, device Out0& x_15, const device In0& x_19, uint3 gl_WorkGroupID_param, thread uint3* const tint_symbol_2) {
|
||||
*(tint_symbol_2) = gl_WorkGroupID_param;
|
||||
main_1(x_13, x_17, x_15, x_19, tint_symbol_2);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint3 gl_WorkGroupID_param [[threadgroup_position_in_grid]], const device In2& x_13 [[buffer(1)]], const device In1& x_17 [[buffer(2)]], device Out0& x_15 [[buffer(0)]], const device In0& x_19 [[buffer(3)]]) {
|
||||
thread uint3 tint_symbol_3 = 0u;
|
||||
tint_symbol_inner(x_13, x_17, x_15, x_19, gl_WorkGroupID_param, &(tint_symbol_3));
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -1,233 +0,0 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 132
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_WorkGroupID_param_1
|
||||
OpExecutionMode %main LocalSize 4 1 1
|
||||
OpName %gl_WorkGroupID_param_1 "gl_WorkGroupID_param_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 %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_inner "main_inner"
|
||||
OpName %gl_WorkGroupID_param "gl_WorkGroupID_param"
|
||||
OpName %main "main"
|
||||
OpDecorate %gl_WorkGroupID_param_1 BuiltIn WorkgroupId
|
||||
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
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%gl_WorkGroupID_param_1 = OpVariable %_ptr_Input_v3uint Input
|
||||
%_ptr_Private_v3uint = OpTypePointer Private %v3uint
|
||||
%7 = OpConstantNull %v3uint
|
||||
%gl_WorkGroupID = OpVariable %_ptr_Private_v3uint Private %7
|
||||
%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
|
||||
%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
|
||||
%123 = OpTypeFunction %void %v3uint
|
||||
%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_inner = OpFunction %void None %123
|
||||
%gl_WorkGroupID_param = OpFunctionParameter %v3uint
|
||||
%126 = OpLabel
|
||||
OpStore %gl_WorkGroupID %gl_WorkGroupID_param
|
||||
%127 = OpFunctionCall %void %main_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %27
|
||||
%129 = OpLabel
|
||||
%131 = OpLoad %v3uint %gl_WorkGroupID_param_1
|
||||
%130 = OpFunctionCall %void %main_inner %131
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -1,120 +0,0 @@
|
||||
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();
|
||||
}
|
||||
@@ -1,72 +0,0 @@
|
||||
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;
|
||||
};
|
||||
|
||||
void main_inner(uint3 gl_WorkGroupID_param) {
|
||||
gl_WorkGroupID = gl_WorkGroupID_param;
|
||||
main_1();
|
||||
}
|
||||
|
||||
[numthreads(4, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
main_inner(tint_symbol.gl_WorkGroupID_param);
|
||||
return;
|
||||
}
|
||||
@@ -1,112 +0,0 @@
|
||||
#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_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;
|
||||
uint const x_58 = (*(tint_symbol_1)).x;
|
||||
base_index_in = (128u * x_58);
|
||||
uint const x_61 = (*(tint_symbol_1)).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 = as_type<int>((as_type<uint>(x_88) - as_type<uint>(1)));
|
||||
int const x_90 = index_in1;
|
||||
index_in1 = as_type<int>((as_type<uint>(x_90) - as_type<uint>(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 = as_type<int>((as_type<uint>(x_103) - as_type<uint>(1)));
|
||||
int const x_105 = index_in1;
|
||||
index_in1 = as_type<int>((as_type<uint>(x_105) - as_type<uint>(1)));
|
||||
}
|
||||
int const x_107 = condition_index;
|
||||
int const x_110 = x_13.data_in2.arr[as_type<int>((as_type<uint>(x_107) + as_type<uint>(1)))];
|
||||
int const x_111 = condition_index;
|
||||
condition_index = as_type<int>((as_type<uint>(x_111) + as_type<uint>(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 = as_type<int>((as_type<uint>(x_119) + as_type<uint>(1)));
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
void tint_symbol_inner(const device In2& x_13, const device In1& x_17, device Out0& x_15, const device In0& x_19, uint3 gl_WorkGroupID_param, thread uint3* const tint_symbol_2) {
|
||||
*(tint_symbol_2) = gl_WorkGroupID_param;
|
||||
main_1(x_13, x_17, x_15, x_19, tint_symbol_2);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(uint3 gl_WorkGroupID_param [[threadgroup_position_in_grid]], const device In2& x_13 [[buffer(1)]], const device In1& x_17 [[buffer(2)]], device Out0& x_15 [[buffer(0)]], const device In0& x_19 [[buffer(3)]]) {
|
||||
thread uint3 tint_symbol_3 = 0u;
|
||||
tint_symbol_inner(x_13, x_17, x_15, x_19, gl_WorkGroupID_param, &(tint_symbol_3));
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -1,233 +0,0 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 132
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_WorkGroupID_param_1
|
||||
OpExecutionMode %main LocalSize 4 1 1
|
||||
OpName %gl_WorkGroupID_param_1 "gl_WorkGroupID_param_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 %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_inner "main_inner"
|
||||
OpName %gl_WorkGroupID_param "gl_WorkGroupID_param"
|
||||
OpName %main "main"
|
||||
OpDecorate %gl_WorkGroupID_param_1 BuiltIn WorkgroupId
|
||||
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
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%gl_WorkGroupID_param_1 = OpVariable %_ptr_Input_v3uint Input
|
||||
%_ptr_Private_v3uint = OpTypePointer Private %v3uint
|
||||
%7 = OpConstantNull %v3uint
|
||||
%gl_WorkGroupID = OpVariable %_ptr_Private_v3uint Private %7
|
||||
%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
|
||||
%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
|
||||
%123 = OpTypeFunction %void %v3uint
|
||||
%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_inner = OpFunction %void None %123
|
||||
%gl_WorkGroupID_param = OpFunctionParameter %v3uint
|
||||
%126 = OpLabel
|
||||
OpStore %gl_WorkGroupID %gl_WorkGroupID_param
|
||||
%127 = OpFunctionCall %void %main_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %27
|
||||
%129 = OpLabel
|
||||
%131 = OpLoad %v3uint %gl_WorkGroupID_param_1
|
||||
%130 = OpFunctionCall %void %main_inner %131
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@@ -1,120 +0,0 @@
|
||||
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