e2e tests: generate missing expected files
Change-Id: I6b4aee4bb08b0f4c02c015f469edc24be1623fc3 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/73340 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: James Price <jrprice@google.com> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
Antonio Maiorano
Tint LUCI CQ
parent
667bc2d929
commit
a9d6c34d86
|
|
@ -0,0 +1,16 @@
|
|||
[numthreads(1, 1, 1)]
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
struct A {
|
||||
int a;
|
||||
};
|
||||
struct B {
|
||||
int b;
|
||||
};
|
||||
|
||||
B f(A a) {
|
||||
const B tint_symbol = (B)0;
|
||||
return tint_symbol;
|
||||
}
|
||||
|
|
@ -0,0 +1,34 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 13
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
|
||||
OpExecutionMode %unused_entry_point LocalSize 1 1 1
|
||||
OpName %unused_entry_point "unused_entry_point"
|
||||
OpName %B "B"
|
||||
OpMemberName %B 0 "b"
|
||||
OpName %A "A"
|
||||
OpMemberName %A 0 "a"
|
||||
OpName %f "f"
|
||||
OpName %a "a"
|
||||
OpMemberDecorate %B 0 Offset 0
|
||||
OpMemberDecorate %A 0 Offset 0
|
||||
%void = OpTypeVoid
|
||||
%1 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%B = OpTypeStruct %int
|
||||
%A = OpTypeStruct %int
|
||||
%5 = OpTypeFunction %B %A
|
||||
%12 = OpConstantNull %B
|
||||
%unused_entry_point = OpFunction %void None %1
|
||||
%4 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %B None %5
|
||||
%a = OpFunctionParameter %A
|
||||
%11 = OpLabel
|
||||
OpReturnValue %12
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
struct A {
|
||||
a : i32;
|
||||
};
|
||||
|
||||
struct B {
|
||||
b : i32;
|
||||
};
|
||||
|
||||
fn f(a : A) -> B {
|
||||
return B();
|
||||
}
|
||||
|
|
@ -0,0 +1,21 @@
|
|||
cbuffer cbuffer_ubo : register(b0, space0) {
|
||||
uint4 ubo[1];
|
||||
};
|
||||
|
||||
struct S {
|
||||
int data[64];
|
||||
};
|
||||
|
||||
RWByteAddressBuffer result : register(u1, space0);
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
S s = (S)0;
|
||||
{
|
||||
int tint_symbol_2[64] = s.data;
|
||||
tint_symbol_2[asint(ubo[0].x)] = 1;
|
||||
s.data = tint_symbol_2;
|
||||
}
|
||||
result.Store(0u, asuint(s.data[3]));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
cbuffer cbuffer_ubo : register(b0, space0) {
|
||||
uint4 ubo[1];
|
||||
};
|
||||
|
||||
struct S {
|
||||
int data[64];
|
||||
};
|
||||
|
||||
RWByteAddressBuffer result : register(u1, space0);
|
||||
|
||||
void x(inout S p) {
|
||||
{
|
||||
int tint_symbol_2[64] = p.data;
|
||||
tint_symbol_2[asint(ubo[0].x)] = 1;
|
||||
p.data = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
S s = (S)0;
|
||||
x(s);
|
||||
result.Store(0u, asuint(s.data[3]));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,21 @@
|
|||
cbuffer cbuffer_ubo : register(b0, space0) {
|
||||
uint4 ubo[1];
|
||||
};
|
||||
|
||||
struct S {
|
||||
int data[64];
|
||||
};
|
||||
|
||||
RWByteAddressBuffer result : register(u1, space0);
|
||||
static S s = (S)0;
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
{
|
||||
int tint_symbol_2[64] = s.data;
|
||||
tint_symbol_2[asint(ubo[0].x)] = 1;
|
||||
s.data = tint_symbol_2;
|
||||
}
|
||||
result.Store(0u, asuint(s.data[3]));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
cbuffer cbuffer_ubo : register(b0, space0) {
|
||||
uint4 ubo[1];
|
||||
};
|
||||
|
||||
struct S {
|
||||
int data[64];
|
||||
};
|
||||
|
||||
RWByteAddressBuffer result : register(u1, space0);
|
||||
static S s = (S)0;
|
||||
|
||||
void x(inout S p) {
|
||||
{
|
||||
int tint_symbol_2[64] = p.data;
|
||||
tint_symbol_2[asint(ubo[0].x)] = 1;
|
||||
p.data = tint_symbol_2;
|
||||
}
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void f() {
|
||||
x(s);
|
||||
result.Store(0u, asuint(s.data[3]));
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,38 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
|
||||
return lhs * vec<T, N>(rhs);
|
||||
}
|
||||
|
||||
template<typename T, int N, int M>
|
||||
inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
|
||||
return vec<T, M>(lhs) * rhs;
|
||||
}
|
||||
|
||||
struct Simulation {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ float3 arr[8];
|
||||
};
|
||||
struct Particle {
|
||||
/* 0x0000 */ tint_array_wrapper position;
|
||||
/* 0x0080 */ float lifetime;
|
||||
/* 0x0084 */ int8_t tint_pad[12];
|
||||
/* 0x0090 */ float4 color;
|
||||
/* 0x00a0 */ packed_float3 velocity;
|
||||
/* 0x00ac */ int8_t tint_pad_1[4];
|
||||
};
|
||||
struct Particles {
|
||||
/* 0x0000 */ Particle p[1];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const device Particles* tint_symbol_1 [[buffer(1)]], const constant Simulation* tint_symbol_2 [[buffer(0)]]) {
|
||||
Particle particle = (*(tint_symbol_1)).p[0];
|
||||
particle.position.arr[(*(tint_symbol_2)).i] = particle.position.arr[(*(tint_symbol_2)).i];
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,78 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 37
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Particles "Particles"
|
||||
OpMemberName %Particles 0 "p"
|
||||
OpName %Particle "Particle"
|
||||
OpMemberName %Particle 0 "position"
|
||||
OpMemberName %Particle 1 "lifetime"
|
||||
OpMemberName %Particle 2 "color"
|
||||
OpMemberName %Particle 3 "velocity"
|
||||
OpName %particles "particles"
|
||||
OpName %Simulation "Simulation"
|
||||
OpMemberName %Simulation 0 "i"
|
||||
OpName %sim "sim"
|
||||
OpName %main "main"
|
||||
OpName %particle "particle"
|
||||
OpDecorate %Particles Block
|
||||
OpMemberDecorate %Particles 0 Offset 0
|
||||
OpMemberDecorate %Particle 0 Offset 0
|
||||
OpDecorate %_arr_v3float_uint_8 ArrayStride 16
|
||||
OpMemberDecorate %Particle 1 Offset 128
|
||||
OpMemberDecorate %Particle 2 Offset 144
|
||||
OpMemberDecorate %Particle 3 Offset 160
|
||||
OpDecorate %_runtimearr_Particle ArrayStride 176
|
||||
OpDecorate %particles NonWritable
|
||||
OpDecorate %particles DescriptorSet 1
|
||||
OpDecorate %particles Binding 3
|
||||
OpDecorate %Simulation Block
|
||||
OpMemberDecorate %Simulation 0 Offset 0
|
||||
OpDecorate %sim NonWritable
|
||||
OpDecorate %sim DescriptorSet 1
|
||||
OpDecorate %sim Binding 4
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_v3float_uint_8 = OpTypeArray %v3float %uint_8
|
||||
%v4float = OpTypeVector %float 4
|
||||
%Particle = OpTypeStruct %_arr_v3float_uint_8 %float %v4float %v3float
|
||||
%_runtimearr_Particle = OpTypeRuntimeArray %Particle
|
||||
%Particles = OpTypeStruct %_runtimearr_Particle
|
||||
%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
|
||||
%particles = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
|
||||
%Simulation = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Simulation = OpTypePointer Uniform %Simulation
|
||||
%sim = OpVariable %_ptr_Uniform_Simulation Uniform
|
||||
%void = OpTypeVoid
|
||||
%15 = OpTypeFunction %void
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%int = OpTypeInt 32 1
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle
|
||||
%_ptr_Function_Particle = OpTypePointer Function %Particle
|
||||
%27 = OpConstantNull %Particle
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%_ptr_Function_v3float = OpTypePointer Function %v3float
|
||||
%main = OpFunction %void None %15
|
||||
%18 = OpLabel
|
||||
%particle = OpVariable %_ptr_Function_Particle Function %27
|
||||
%23 = OpAccessChain %_ptr_StorageBuffer_Particle %particles %uint_0 %int_0
|
||||
%24 = OpLoad %Particle %23
|
||||
OpStore %particle %24
|
||||
%29 = OpAccessChain %_ptr_Uniform_uint %sim %uint_0
|
||||
%30 = OpLoad %uint %29
|
||||
%32 = OpAccessChain %_ptr_Function_v3float %particle %uint_0 %30
|
||||
%33 = OpAccessChain %_ptr_Uniform_uint %sim %uint_0
|
||||
%34 = OpLoad %uint %33
|
||||
%35 = OpAccessChain %_ptr_Function_v3float %particle %uint_0 %34
|
||||
%36 = OpLoad %v3float %35
|
||||
OpStore %32 %36
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,24 @@
|
|||
struct Simulation {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct Particle {
|
||||
position : array<vec3<f32>, 8>;
|
||||
lifetime : f32;
|
||||
color : vec4<f32>;
|
||||
velocity : vec3<f32>;
|
||||
};
|
||||
|
||||
struct Particles {
|
||||
p : array<Particle>;
|
||||
};
|
||||
|
||||
[[group(1), binding(3)]] var<storage, read> particles : Particles;
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> sim : Simulation;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var particle = particles.p[0];
|
||||
particle.position[sim.i] = particle.position[sim.i];
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
|
|
@ -0,0 +1,21 @@
|
|||
cbuffer cbuffer_constants : register(b0, space1) {
|
||||
uint4 constants[1];
|
||||
};
|
||||
|
||||
RWByteAddressBuffer result : register(u1, space1);
|
||||
|
||||
struct S {
|
||||
uint data[3];
|
||||
};
|
||||
|
||||
static S s = (S)0;
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main() {
|
||||
{
|
||||
uint tint_symbol_1[3] = s.data;
|
||||
tint_symbol_1[constants[0].x] = 0u;
|
||||
s.data = tint_symbol_1;
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
|
||||
s1.a1.arr[(*(tint_symbol_1)).i] = v;
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,84 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 44
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpName %i "i"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%19 = OpConstantNull %OuterS
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%23 = OpConstantNull %int
|
||||
%int_4 = OpConstant %int 4
|
||||
%bool = OpTypeBool
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%int_1 = OpConstant %int 1
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %19
|
||||
%i = OpVariable %_ptr_Function_int Function %23
|
||||
OpStore %i %int_0
|
||||
OpBranch %24
|
||||
%24 = OpLabel
|
||||
OpLoopMerge %25 %26 None
|
||||
OpBranch %27
|
||||
%27 = OpLabel
|
||||
%29 = OpLoad %int %i
|
||||
%31 = OpSLessThan %bool %29 %int_4
|
||||
%28 = OpLogicalNot %bool %31
|
||||
OpSelectionMerge %33 None
|
||||
OpBranchConditional %28 %34 %33
|
||||
%34 = OpLabel
|
||||
OpBranch %25
|
||||
%33 = OpLabel
|
||||
%37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%38 = OpLoad %uint %37
|
||||
%39 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %38
|
||||
%40 = OpLoad %InnerS %v
|
||||
OpStore %39 %40
|
||||
OpBranch %26
|
||||
%26 = OpLabel
|
||||
%41 = OpLoad %int %i
|
||||
%43 = OpIAdd %int %41 %int_1
|
||||
OpStore %i %43
|
||||
OpBranch %24
|
||||
%25 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,22 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
for(var i : i32 = 0; (i < 4); i = (i + 1)) {
|
||||
s1.a1[uniforms.i] = v;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
for(int i = 0; (i < 4); s1.a1.arr[(*(tint_symbol_1)).i] = v) {
|
||||
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,84 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 44
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpName %i "i"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%19 = OpConstantNull %OuterS
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%23 = OpConstantNull %int
|
||||
%int_4 = OpConstant %int 4
|
||||
%bool = OpTypeBool
|
||||
%int_1 = OpConstant %int 1
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %19
|
||||
%i = OpVariable %_ptr_Function_int Function %23
|
||||
OpStore %i %int_0
|
||||
OpBranch %24
|
||||
%24 = OpLabel
|
||||
OpLoopMerge %25 %26 None
|
||||
OpBranch %27
|
||||
%27 = OpLabel
|
||||
%29 = OpLoad %int %i
|
||||
%31 = OpSLessThan %bool %29 %int_4
|
||||
%28 = OpLogicalNot %bool %31
|
||||
OpSelectionMerge %33 None
|
||||
OpBranchConditional %28 %34 %33
|
||||
%34 = OpLabel
|
||||
OpBranch %25
|
||||
%33 = OpLabel
|
||||
%35 = OpLoad %int %i
|
||||
%37 = OpIAdd %int %35 %int_1
|
||||
OpStore %i %37
|
||||
OpBranch %26
|
||||
%26 = OpLabel
|
||||
%40 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%41 = OpLoad %uint %40
|
||||
%42 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %41
|
||||
%43 = OpLoad %InnerS %v
|
||||
OpStore %42 %43
|
||||
OpBranch %24
|
||||
%25 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,22 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
for(var i : i32 = 0; (i < 4); s1.a1[uniforms.i] = v) {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
int i = 0;
|
||||
for(s1.a1.arr[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,84 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 44
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpName %i "i"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%19 = OpConstantNull %OuterS
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%23 = OpConstantNull %int
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%int_4 = OpConstant %int 4
|
||||
%bool = OpTypeBool
|
||||
%int_1 = OpConstant %int 1
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %19
|
||||
%i = OpVariable %_ptr_Function_int Function %23
|
||||
OpStore %i %int_0
|
||||
%26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%27 = OpLoad %uint %26
|
||||
%28 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %27
|
||||
%29 = OpLoad %InnerS %v
|
||||
OpStore %28 %29
|
||||
OpBranch %30
|
||||
%30 = OpLabel
|
||||
OpLoopMerge %31 %32 None
|
||||
OpBranch %33
|
||||
%33 = OpLabel
|
||||
%35 = OpLoad %int %i
|
||||
%37 = OpSLessThan %bool %35 %int_4
|
||||
%34 = OpLogicalNot %bool %37
|
||||
OpSelectionMerge %39 None
|
||||
OpBranchConditional %34 %40 %39
|
||||
%40 = OpLabel
|
||||
OpBranch %31
|
||||
%39 = OpLabel
|
||||
OpBranch %32
|
||||
%32 = OpLabel
|
||||
%41 = OpLoad %int %i
|
||||
%43 = OpIAdd %int %41 %int_1
|
||||
OpStore %i %43
|
||||
OpBranch %30
|
||||
%31 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,22 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
var i : i32 = 0;
|
||||
for(s1.a1[uniforms.i] = v; (i < 4); i = (i + 1)) {
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,36 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
/* 0x0004 */ uint j;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct S1 {
|
||||
tint_array_wrapper a2;
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
S1 arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper_1 a1;
|
||||
};
|
||||
|
||||
uint getNextIndex(thread uint* const tint_symbol_1) {
|
||||
*(tint_symbol_1) = (*(tint_symbol_1) + 1u);
|
||||
return *(tint_symbol_1);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_3 [[buffer(0)]]) {
|
||||
thread uint tint_symbol_2 = 0u;
|
||||
InnerS v = {};
|
||||
OuterS s = {};
|
||||
s.a1.arr[getNextIndex(&(tint_symbol_2))].a2.arr[(*(tint_symbol_3)).j] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,79 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 39
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %nextIndex "nextIndex"
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpMemberName %Uniforms 1 "j"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %getNextIndex "getNextIndex"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %S1 "S1"
|
||||
OpMemberName %S1 0 "a2"
|
||||
OpName %s "s"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpMemberDecorate %Uniforms 1 Offset 4
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %S1 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
OpDecorate %_arr_S1_uint_8 ArrayStride 32
|
||||
%uint = OpTypeInt 32 0
|
||||
%_ptr_Private_uint = OpTypePointer Private %uint
|
||||
%4 = OpConstantNull %uint
|
||||
%nextIndex = OpVariable %_ptr_Private_uint Private %4
|
||||
%Uniforms = OpTypeStruct %uint %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%8 = OpTypeFunction %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%void = OpTypeVoid
|
||||
%15 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%23 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%S1 = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_S1_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%31 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%getNextIndex = OpFunction %uint None %8
|
||||
%10 = OpLabel
|
||||
%11 = OpLoad %uint %nextIndex
|
||||
%13 = OpIAdd %uint %11 %uint_1
|
||||
OpStore %nextIndex %13
|
||||
%14 = OpLoad %uint %nextIndex
|
||||
OpReturnValue %14
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %15
|
||||
%18 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %23
|
||||
%s = OpVariable %_ptr_Function_OuterS Function %31
|
||||
%33 = OpFunctionCall %uint %getNextIndex
|
||||
%35 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%36 = OpLoad %uint %35
|
||||
%37 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %33 %uint_0 %36
|
||||
%38 = OpLoad %InnerS %v
|
||||
OpStore %37 %38
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,32 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
j : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct S1 {
|
||||
a2 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<S1, 8>;
|
||||
};
|
||||
|
||||
var<private> nextIndex : u32;
|
||||
|
||||
fn getNextIndex() -> u32 {
|
||||
nextIndex = (nextIndex + 1u);
|
||||
return nextIndex;
|
||||
}
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s : OuterS;
|
||||
s.a1[getNextIndex()].a2[uniforms.j] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,23 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
s1.a1.arr[(*(tint_symbol_1)).i] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,55 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 26
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%19 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %19
|
||||
%22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%23 = OpLoad %uint %22
|
||||
%24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23
|
||||
%25 = OpLoad %InnerS %v
|
||||
OpStore %24 %25
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,20 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
s1.a1[uniforms.i] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,27 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
/* 0x0004 */ uint j;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
tint_array_wrapper_1 arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
s1.a1.arr[(*(tint_symbol_1)).i].arr[(*(tint_symbol_1)).j] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,62 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 30
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpMemberName %Uniforms 1 "j"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpMemberDecorate %Uniforms 1 Offset 4
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
OpDecorate %_arr__arr_InnerS_uint_8_uint_8 ArrayStride 32
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%_arr__arr_InnerS_uint_8_uint_8 = OpTypeArray %_arr_InnerS_uint_8 %uint_8
|
||||
%OuterS = OpTypeStruct %_arr__arr_InnerS_uint_8_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%20 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %20
|
||||
%23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%24 = OpLoad %uint %23
|
||||
%26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%27 = OpLoad %uint %26
|
||||
%28 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %24 %27
|
||||
%29 = OpLoad %InnerS %v
|
||||
OpStore %28 %29
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,21 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
j : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<array<InnerS, 8>, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
s1.a1[uniforms.i][uniforms.j] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,26 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct S1 {
|
||||
InnerS s2;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
S1 arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
s1.a1.arr[(*(tint_symbol_1)).i].s2 = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,59 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 27
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %S1 "S1"
|
||||
OpMemberName %S1 0 "s2"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %S1 0 Offset 0
|
||||
OpDecorate %_arr_S1_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%S1 = OpTypeStruct %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_S1_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%20 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %20
|
||||
%23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%24 = OpLoad %uint %23
|
||||
%25 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %24 %uint_0
|
||||
%26 = OpLoad %InnerS %v
|
||||
OpStore %25 %26
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,24 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct S1 {
|
||||
s2 : InnerS;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<S1, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
s1.a1[uniforms.i].s2 = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,30 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
/* 0x0004 */ uint j;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct S1 {
|
||||
tint_array_wrapper a2;
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
S1 arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper_1 a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s = {};
|
||||
s.a1.arr[(*(tint_symbol_1)).i].a2.arr[(*(tint_symbol_1)).j] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,66 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 31
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpMemberName %Uniforms 1 "j"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %S1 "S1"
|
||||
OpMemberName %S1 0 "a2"
|
||||
OpName %s "s"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpMemberDecorate %Uniforms 1 Offset 4
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %S1 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
OpDecorate %_arr_S1_uint_8 ArrayStride 32
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%S1 = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_S1_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%21 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s = OpVariable %_ptr_Function_OuterS Function %21
|
||||
%24 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%25 = OpLoad %uint %24
|
||||
%27 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%28 = OpLoad %uint %27
|
||||
%29 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %25 %uint_0 %28
|
||||
%30 = OpLoad %InnerS %v
|
||||
OpStore %29 %30
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
j : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct S1 {
|
||||
a2 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<S1, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s : OuterS;
|
||||
s.a1[uniforms.i].a2[uniforms.j] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,19 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
/* 0x0000 */ int v;
|
||||
};
|
||||
struct OuterS {
|
||||
/* 0x0000 */ InnerS a1[1];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(device OuterS* tint_symbol_1 [[buffer(1)]], const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
(*(tint_symbol_1)).a1[(*(tint_symbol_2)).i] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,57 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 25
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %s1 "s1"
|
||||
OpName %main "main"
|
||||
OpName %v "v"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpDecorate %OuterS Block
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpDecorate %_runtimearr_InnerS ArrayStride 4
|
||||
OpDecorate %s1 Binding 0
|
||||
OpDecorate %s1 DescriptorSet 0
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_runtimearr_InnerS = OpTypeRuntimeArray %InnerS
|
||||
%OuterS = OpTypeStruct %_runtimearr_InnerS
|
||||
%_ptr_StorageBuffer_OuterS = OpTypePointer StorageBuffer %OuterS
|
||||
%s1 = OpVariable %_ptr_StorageBuffer_OuterS StorageBuffer
|
||||
%void = OpTypeVoid
|
||||
%11 = OpTypeFunction %void
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%17 = OpConstantNull %InnerS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%_ptr_StorageBuffer_InnerS = OpTypePointer StorageBuffer %InnerS
|
||||
%main = OpFunction %void None %11
|
||||
%14 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %17
|
||||
%20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%21 = OpLoad %uint %20
|
||||
%23 = OpAccessChain %_ptr_StorageBuffer_InnerS %s1 %uint_0 %21
|
||||
%24 = OpLoad %InnerS %v
|
||||
OpStore %23 %24
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,21 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[binding(0), group(0)]] var<storage, read_write> s1 : OuterS;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
s1.a1[uniforms.i] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,26 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
/* 0x0004 */ uint j;
|
||||
};
|
||||
struct InnerS {
|
||||
/* 0x0000 */ int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ InnerS arr[8];
|
||||
};
|
||||
struct S1 {
|
||||
/* 0x0000 */ tint_array_wrapper a2;
|
||||
};
|
||||
struct OuterS {
|
||||
/* 0x0000 */ S1 a1[1];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(device OuterS* tint_symbol_1 [[buffer(1)]], const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
(*(tint_symbol_1)).a1[(*(tint_symbol_2)).i].a2.arr[(*(tint_symbol_2)).j] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,69 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 31
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpMemberName %Uniforms 1 "j"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %S1 "S1"
|
||||
OpMemberName %S1 0 "a2"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %s "s"
|
||||
OpName %main "main"
|
||||
OpName %v "v"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpMemberDecorate %Uniforms 1 Offset 4
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpDecorate %OuterS Block
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %S1 0 Offset 0
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
OpDecorate %_runtimearr_S1 ArrayStride 32
|
||||
OpDecorate %s Binding 0
|
||||
OpDecorate %s DescriptorSet 0
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%S1 = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_runtimearr_S1 = OpTypeRuntimeArray %S1
|
||||
%OuterS = OpTypeStruct %_runtimearr_S1
|
||||
%_ptr_StorageBuffer_OuterS = OpTypePointer StorageBuffer %OuterS
|
||||
%s = OpVariable %_ptr_StorageBuffer_OuterS StorageBuffer
|
||||
%void = OpTypeVoid
|
||||
%14 = OpTypeFunction %void
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%20 = OpConstantNull %InnerS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_StorageBuffer_InnerS = OpTypePointer StorageBuffer %InnerS
|
||||
%main = OpFunction %void None %14
|
||||
%17 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %20
|
||||
%23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%24 = OpLoad %uint %23
|
||||
%26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%27 = OpLoad %uint %26
|
||||
%29 = OpAccessChain %_ptr_StorageBuffer_InnerS %s %uint_0 %24 %uint_0 %27
|
||||
%30 = OpLoad %InnerS %v
|
||||
OpStore %29 %30
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,26 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
j : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct S1 {
|
||||
a2 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<S1>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[binding(0), group(0)]] var<storage, read_write> s : OuterS;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
s.a1[uniforms.i].a2[uniforms.j] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,17 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct OuterS {
|
||||
float2x4 m1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
OuterS s1 = {};
|
||||
s1.m1[(*(tint_symbol_1)).i] = float4(1.0f);
|
||||
s1.m1[(*(tint_symbol_1)).i][(*(tint_symbol_1)).i] = 1.0f;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,57 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 30
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "m1"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 ColMajor
|
||||
OpMemberDecorate %OuterS 0 MatrixStride 16
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%mat2v4float = OpTypeMatrix %v4float 2
|
||||
%OuterS = OpTypeStruct %mat2v4float
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%15 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%_ptr_Function_v4float = OpTypePointer Function %v4float
|
||||
%float_1 = OpConstant %float 1
|
||||
%23 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %15
|
||||
%18 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%19 = OpLoad %uint %18
|
||||
%21 = OpAccessChain %_ptr_Function_v4float %s1 %uint_0 %19
|
||||
OpStore %21 %23
|
||||
%24 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%25 = OpLoad %uint %24
|
||||
%26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%27 = OpLoad %uint %26
|
||||
%29 = OpAccessChain %_ptr_Function_float %s1 %uint_0 %25 %27
|
||||
OpStore %29 %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,16 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
m1 : mat2x4<f32>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var s1 : OuterS;
|
||||
s1.m1[uniforms.i] = vec4<f32>(1.0);
|
||||
s1.m1[uniforms.i][uniforms.i] = 1.0;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
tint_array_wrapper a2;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
s1.a1.arr[(*(tint_symbol_1)).i] = v;
|
||||
s1.a2.arr[(*(tint_symbol_1)).i] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,63 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 31
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpMemberName %OuterS 1 "a2"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
OpMemberDecorate %OuterS 1 Offset 32
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8 %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%19 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %19
|
||||
%22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%23 = OpLoad %uint %22
|
||||
%24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23
|
||||
%25 = OpLoad %InnerS %v
|
||||
OpStore %24 %25
|
||||
%27 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%28 = OpLoad %uint %27
|
||||
%29 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_1 %28
|
||||
%30 = OpLoad %InnerS %v
|
||||
OpStore %29 %30
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,22 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
a2 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
s1.a1[uniforms.i] = v;
|
||||
s1.a2[uniforms.i] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,26 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct S1 {
|
||||
tint_array_wrapper a;
|
||||
};
|
||||
struct OuterS {
|
||||
S1 s2;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
s1.s2.a.arr[(*(tint_symbol_1)).i] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,59 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 27
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "s2"
|
||||
OpName %S1 "S1"
|
||||
OpMemberName %S1 0 "a"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %S1 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%S1 = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%OuterS = OpTypeStruct %S1
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%20 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %20
|
||||
%23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%24 = OpLoad %uint %23
|
||||
%25 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %uint_0 %24
|
||||
%26 = OpLoad %InnerS %v
|
||||
OpStore %25 %26
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,24 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct S1 {
|
||||
a : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
s2 : S1;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
s1.s2.a[uniforms.i] = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,16 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct OuterS {
|
||||
float3 v1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
OuterS s1 = {};
|
||||
s1.v1[(*(tint_symbol_1)).i] = 1.0f;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,46 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 22
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "v1"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%OuterS = OpTypeStruct %v3float
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%14 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%float_1 = OpConstant %float 1
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %14
|
||||
%17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%18 = OpLoad %uint %17
|
||||
%20 = OpAccessChain %_ptr_Function_float %s1 %uint_0 %18
|
||||
OpStore %20 %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
v1 : vec3<f32>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var s1 : OuterS;
|
||||
s1.v1[uniforms.i] = 1.0;
|
||||
}
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
uint arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
uint f(uint i) {
|
||||
return (i + 1u);
|
||||
}
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
OuterS s1 = {};
|
||||
float3 v = 0.0f;
|
||||
v[s1.a1.arr[(*(tint_symbol_1)).i]] = 1.0f;
|
||||
v[f(s1.a1.arr[(*(tint_symbol_1)).i])] = 1.0f;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,73 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 42
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %f "f"
|
||||
OpName %i "i"
|
||||
OpName %main "main"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpName %v "v"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_uint_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%5 = OpTypeFunction %uint %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%void = OpTypeVoid
|
||||
%11 = OpTypeFunction %void
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_uint_uint_8 = OpTypeArray %uint %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_uint_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%20 = OpConstantNull %OuterS
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%_ptr_Function_v3float = OpTypePointer Function %v3float
|
||||
%25 = OpConstantNull %v3float
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%float_1 = OpConstant %float 1
|
||||
%f = OpFunction %uint None %5
|
||||
%i = OpFunctionParameter %uint
|
||||
%8 = OpLabel
|
||||
%10 = OpIAdd %uint %i %uint_1
|
||||
OpReturnValue %10
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %11
|
||||
%14 = OpLabel
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %20
|
||||
%v = OpVariable %_ptr_Function_v3float Function %25
|
||||
%28 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%29 = OpLoad %uint %28
|
||||
%31 = OpAccessChain %_ptr_Function_uint %s1 %uint_0 %29
|
||||
%32 = OpLoad %uint %31
|
||||
%34 = OpAccessChain %_ptr_Function_float %v %32
|
||||
OpStore %34 %float_1
|
||||
%37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%38 = OpLoad %uint %37
|
||||
%39 = OpAccessChain %_ptr_Function_uint %s1 %uint_0 %38
|
||||
%40 = OpLoad %uint %39
|
||||
%36 = OpFunctionCall %uint %f %40
|
||||
%41 = OpAccessChain %_ptr_Function_float %v %36
|
||||
OpStore %41 %float_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,21 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<u32, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
fn f(i : u32) -> u32 {
|
||||
return (i + 1u);
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var s1 : OuterS;
|
||||
var v : vec3<f32>;
|
||||
v[s1.a1[uniforms.i]] = 1.0;
|
||||
v[f(s1.a1[uniforms.i])] = 1.0;
|
||||
}
|
||||
|
|
@ -0,0 +1,24 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
|
||||
InnerS v = {};
|
||||
OuterS s1 = {};
|
||||
uint const p_save = (*(tint_symbol_1)).i;
|
||||
s1.a1.arr[p_save] = v;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,55 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 26
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %main "main"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %v "v"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%13 = OpConstantNull %InnerS
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%19 = OpConstantNull %OuterS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %13
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %19
|
||||
%22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%23 = OpLoad %uint %22
|
||||
%24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23
|
||||
%25 = OpLoad %InnerS %v
|
||||
OpStore %24 %25
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,21 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var v : InnerS;
|
||||
var s1 : OuterS;
|
||||
let p = &(s1.a1[uniforms.i]);
|
||||
*(p) = v;
|
||||
}
|
||||
|
|
@ -0,0 +1,27 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint i;
|
||||
};
|
||||
struct InnerS {
|
||||
int v;
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
InnerS arr[8];
|
||||
};
|
||||
struct OuterS {
|
||||
tint_array_wrapper a1;
|
||||
};
|
||||
|
||||
void f(thread OuterS* const p, const constant Uniforms* const tint_symbol_1) {
|
||||
InnerS v = {};
|
||||
(*(p)).a1.arr[(*(tint_symbol_1)).i] = v;
|
||||
}
|
||||
|
||||
kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
|
||||
OuterS s1 = {};
|
||||
f(&(s1), tint_symbol_2);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,64 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 33
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "i"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %OuterS "OuterS"
|
||||
OpMemberName %OuterS 0 "a1"
|
||||
OpName %InnerS "InnerS"
|
||||
OpMemberName %InnerS 0 "v"
|
||||
OpName %f "f"
|
||||
OpName %p "p"
|
||||
OpName %v "v"
|
||||
OpName %main "main"
|
||||
OpName %s1 "s1"
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 1
|
||||
OpDecorate %uniforms Binding 4
|
||||
OpMemberDecorate %OuterS 0 Offset 0
|
||||
OpMemberDecorate %InnerS 0 Offset 0
|
||||
OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
|
||||
%uint = OpTypeInt 32 0
|
||||
%Uniforms = OpTypeStruct %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%void = OpTypeVoid
|
||||
%int = OpTypeInt 32 1
|
||||
%InnerS = OpTypeStruct %int
|
||||
%uint_8 = OpConstant %uint 8
|
||||
%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
|
||||
%OuterS = OpTypeStruct %_arr_InnerS_uint_8
|
||||
%_ptr_Function_OuterS = OpTypePointer Function %OuterS
|
||||
%5 = OpTypeFunction %void %_ptr_Function_OuterS
|
||||
%_ptr_Function_InnerS = OpTypePointer Function %InnerS
|
||||
%18 = OpConstantNull %InnerS
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%26 = OpTypeFunction %void
|
||||
%30 = OpConstantNull %OuterS
|
||||
%f = OpFunction %void None %5
|
||||
%p = OpFunctionParameter %_ptr_Function_OuterS
|
||||
%15 = OpLabel
|
||||
%v = OpVariable %_ptr_Function_InnerS Function %18
|
||||
%22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%23 = OpLoad %uint %22
|
||||
%24 = OpAccessChain %_ptr_Function_InnerS %p %uint_0 %23
|
||||
%25 = OpLoad %InnerS %v
|
||||
OpStore %24 %25
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %26
|
||||
%28 = OpLabel
|
||||
%s1 = OpVariable %_ptr_Function_OuterS Function %30
|
||||
%31 = OpFunctionCall %void %f %s1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,24 @@
|
|||
struct Uniforms {
|
||||
i : u32;
|
||||
};
|
||||
|
||||
struct InnerS {
|
||||
v : i32;
|
||||
};
|
||||
|
||||
struct OuterS {
|
||||
a1 : array<InnerS, 8>;
|
||||
};
|
||||
|
||||
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
fn f(p : ptr<function, OuterS>) {
|
||||
var v : InnerS;
|
||||
(*(p)).a1[uniforms.i] = v;
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(1)]]
|
||||
fn main() {
|
||||
var s1 : OuterS;
|
||||
f(&(s1));
|
||||
}
|
||||
Loading…
Reference in New Issue