test: Add tests to verify variables are zero initialized
Bug: tint:938 Fixed: tint:759 Change-Id: I5794d7f40ec154c55240b163316ed45c7f14b190 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56547 Reviewed-by: James Price <jrprice@google.com> Reviewed-by: Corentin Wallez <cwallez@chromium.org> Commit-Queue: Ben Clayton <bclayton@google.com> Auto-Submit: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
71a1f58537
commit
fb4e751258
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : array<i32, 3>;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,10 @@
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
int arr[3];
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
tint_array_wrapper v = {{0, 0, 0}};
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,13 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
int arr[3];
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
tint_array_wrapper v = {};
|
||||||
|
(void) v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,26 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 14
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %v "v"
|
||||||
|
OpDecorate %_arr_int_uint_3 ArrayStride 4
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%1 = OpTypeFunction %void
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%uint_3 = OpConstant %uint 3
|
||||||
|
%_arr_int_uint_3 = OpTypeArray %int %uint_3
|
||||||
|
%_ptr_Function__arr_int_uint_3 = OpTypePointer Function %_arr_int_uint_3
|
||||||
|
%11 = OpConstantNull %_arr_int_uint_3
|
||||||
|
%main = OpFunction %void None %1
|
||||||
|
%4 = OpLabel
|
||||||
|
%v = OpVariable %_ptr_Function__arr_int_uint_3 Function %11
|
||||||
|
%13 = OpLoad %_arr_int_uint_3 %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : array<i32, 3>;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : mat2x3<f32>;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
float2x3 v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,9 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
float2x3 v = float2x3(0.0f);
|
||||||
|
(void) v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,24 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 13
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %v "v"
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%1 = OpTypeFunction %void
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%v3float = OpTypeVector %float 3
|
||||||
|
%mat2v3float = OpTypeMatrix %v3float 2
|
||||||
|
%_ptr_Function_mat2v3float = OpTypePointer Function %mat2v3float
|
||||||
|
%10 = OpConstantNull %mat2v3float
|
||||||
|
%main = OpFunction %void None %1
|
||||||
|
%4 = OpLabel
|
||||||
|
%v = OpVariable %_ptr_Function_mat2v3float Function %10
|
||||||
|
%12 = OpLoad %mat2v3float %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : mat2x3<f32>;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : i32;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
int v = 0;
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,9 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
int v = 0;
|
||||||
|
(void) v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,22 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 11
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %v "v"
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%1 = OpTypeFunction %void
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%_ptr_Function_int = OpTypePointer Function %int
|
||||||
|
%8 = OpConstantNull %int
|
||||||
|
%main = OpFunction %void None %1
|
||||||
|
%4 = OpLabel
|
||||||
|
%v = OpVariable %_ptr_Function_int Function %8
|
||||||
|
%10 = OpLoad %int %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : i32;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,10 @@
|
||||||
|
struct S {
|
||||||
|
a : i32;
|
||||||
|
b : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : S;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,11 @@
|
||||||
|
struct S {
|
||||||
|
int a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
S v = {0, 0.0f};
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,14 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct S {
|
||||||
|
int a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
S v = {};
|
||||||
|
(void) v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,29 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 13
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %S "S"
|
||||||
|
OpMemberName %S 0 "a"
|
||||||
|
OpMemberName %S 1 "b"
|
||||||
|
OpName %v "v"
|
||||||
|
OpMemberDecorate %S 0 Offset 0
|
||||||
|
OpMemberDecorate %S 1 Offset 4
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%1 = OpTypeFunction %void
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%S = OpTypeStruct %int %float
|
||||||
|
%_ptr_Function_S = OpTypePointer Function %S
|
||||||
|
%10 = OpConstantNull %S
|
||||||
|
%main = OpFunction %void None %1
|
||||||
|
%4 = OpLabel
|
||||||
|
%v = OpVariable %_ptr_Function_S Function %10
|
||||||
|
%12 = OpLoad %S %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,10 @@
|
||||||
|
struct S {
|
||||||
|
a : i32;
|
||||||
|
b : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : S;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : vec3<i32>;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
int3 v = int3(0, 0, 0);
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,9 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
int3 v = 0;
|
||||||
|
(void) v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,23 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 12
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %v "v"
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%1 = OpTypeFunction %void
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%v3int = OpTypeVector %int 3
|
||||||
|
%_ptr_Function_v3int = OpTypePointer Function %v3int
|
||||||
|
%9 = OpConstantNull %v3int
|
||||||
|
%main = OpFunction %void None %1
|
||||||
|
%4 = OpLabel
|
||||||
|
%v = OpVariable %_ptr_Function_v3int Function %9
|
||||||
|
%11 = OpLoad %v3int %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,5 @@
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
var v : vec3<i32>;
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : array<i32, 3>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,11 @@
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
int arr[3];
|
||||||
|
};
|
||||||
|
|
||||||
|
static tint_array_wrapper v = {{0, 0, 0}};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,13 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
int arr[3];
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
thread tint_array_wrapper tint_symbol_1 = {};
|
||||||
|
(void) tint_symbol_1;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,26 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 14
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %_arr_int_uint_3 ArrayStride 4
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%uint_3 = OpConstant %uint 3
|
||||||
|
%_arr_int_uint_3 = OpTypeArray %int %uint_3
|
||||||
|
%_ptr_Private__arr_int_uint_3 = OpTypePointer Private %_arr_int_uint_3
|
||||||
|
%7 = OpConstantNull %_arr_int_uint_3
|
||||||
|
%v = OpVariable %_ptr_Private__arr_int_uint_3 Private %7
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%8 = OpTypeFunction %void
|
||||||
|
%main = OpFunction %void None %8
|
||||||
|
%11 = OpLabel
|
||||||
|
%13 = OpLoad %_arr_int_uint_3 %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : array<i32, 3>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : mat2x3<f32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,7 @@
|
||||||
|
static float2x3 v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,9 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
thread float2x3 tint_symbol_1 = float2x3(0.0f);
|
||||||
|
(void) tint_symbol_1;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,24 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 13
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %main "main"
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%v3float = OpTypeVector %float 3
|
||||||
|
%mat2v3float = OpTypeMatrix %v3float 2
|
||||||
|
%_ptr_Private_mat2v3float = OpTypePointer Private %mat2v3float
|
||||||
|
%6 = OpConstantNull %mat2v3float
|
||||||
|
%v = OpVariable %_ptr_Private_mat2v3float Private %6
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%7 = OpTypeFunction %void
|
||||||
|
%main = OpFunction %void None %7
|
||||||
|
%10 = OpLabel
|
||||||
|
%12 = OpLoad %mat2v3float %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : mat2x3<f32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : i32;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,7 @@
|
||||||
|
static int v = 0;
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,9 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
thread int tint_symbol_1 = 0;
|
||||||
|
(void) tint_symbol_1;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,22 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 11
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %main "main"
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%_ptr_Private_int = OpTypePointer Private %int
|
||||||
|
%4 = OpConstantNull %int
|
||||||
|
%v = OpVariable %_ptr_Private_int Private %4
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%5 = OpTypeFunction %void
|
||||||
|
%main = OpFunction %void None %5
|
||||||
|
%8 = OpLabel
|
||||||
|
%10 = OpLoad %int %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : i32;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,11 @@
|
||||||
|
struct S {
|
||||||
|
a : i32;
|
||||||
|
b : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
var<private> v : S;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,12 @@
|
||||||
|
struct S {
|
||||||
|
int a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
static S v = {0, 0.0f};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,14 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct S {
|
||||||
|
int a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
thread S tint_symbol_1 = {};
|
||||||
|
(void) tint_symbol_1;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,29 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 13
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %S "S"
|
||||||
|
OpMemberName %S 0 "a"
|
||||||
|
OpMemberName %S 1 "b"
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %main "main"
|
||||||
|
OpMemberDecorate %S 0 Offset 0
|
||||||
|
OpMemberDecorate %S 1 Offset 4
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%S = OpTypeStruct %int %float
|
||||||
|
%_ptr_Private_S = OpTypePointer Private %S
|
||||||
|
%6 = OpConstantNull %S
|
||||||
|
%v = OpVariable %_ptr_Private_S Private %6
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%7 = OpTypeFunction %void
|
||||||
|
%main = OpFunction %void None %7
|
||||||
|
%10 = OpLabel
|
||||||
|
%12 = OpLoad %S %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,11 @@
|
||||||
|
struct S {
|
||||||
|
a : i32;
|
||||||
|
b : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
var<private> v : S;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : vec3<i32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,7 @@
|
||||||
|
static int3 v = int3(0, 0, 0);
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main() {
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,9 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol() {
|
||||||
|
thread int3 tint_symbol_1 = 0;
|
||||||
|
(void) tint_symbol_1;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,23 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 12
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main"
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %main "main"
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%v3int = OpTypeVector %int 3
|
||||||
|
%_ptr_Private_v3int = OpTypePointer Private %v3int
|
||||||
|
%5 = OpConstantNull %v3int
|
||||||
|
%v = OpVariable %_ptr_Private_v3int Private %5
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%6 = OpTypeFunction %void
|
||||||
|
%main = OpFunction %void None %6
|
||||||
|
%9 = OpLabel
|
||||||
|
%11 = OpLoad %v3int %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<private> v : vec3<i32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : array<i32, 3>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,21 @@
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
int arr[3];
|
||||||
|
};
|
||||||
|
|
||||||
|
groupshared tint_array_wrapper v;
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint local_invocation_index : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
const tint_array_wrapper tint_symbol_2 = {{0, 0, 0}};
|
||||||
|
v = tint_symbol_2;
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,18 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct tint_array_wrapper {
|
||||||
|
int arr[3];
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||||
|
threadgroup tint_array_wrapper tint_symbol_3;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
tint_array_wrapper const tint_symbol_2 = {.arr={}};
|
||||||
|
tint_symbol_3 = tint_symbol_2;
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
(void) tint_symbol_3;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,43 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 25
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %_arr_int_uint_3 ArrayStride 4
|
||||||
|
OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%uint_3 = OpConstant %uint 3
|
||||||
|
%_arr_int_uint_3 = OpTypeArray %int %uint_3
|
||||||
|
%_ptr_Workgroup__arr_int_uint_3 = OpTypePointer Workgroup %_arr_int_uint_3
|
||||||
|
%v = OpVariable %_ptr_Workgroup__arr_int_uint_3 Workgroup
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_uint Input
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%9 = OpTypeFunction %void
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%19 = OpConstantNull %_arr_int_uint_3
|
||||||
|
%uint_2 = OpConstant %uint 2
|
||||||
|
%uint_264 = OpConstant %uint 264
|
||||||
|
%main = OpFunction %void None %9
|
||||||
|
%12 = OpLabel
|
||||||
|
%13 = OpLoad %uint %tint_symbol
|
||||||
|
%15 = OpIEqual %bool %13 %uint_0
|
||||||
|
OpSelectionMerge %17 None
|
||||||
|
OpBranchConditional %15 %18 %17
|
||||||
|
%18 = OpLabel
|
||||||
|
OpStore %v %19
|
||||||
|
OpBranch %17
|
||||||
|
%17 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
%24 = OpLoad %_arr_int_uint_3 %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : array<i32, 3>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : mat2x3<f32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,16 @@
|
||||||
|
groupshared float2x3 v;
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint local_invocation_index : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,192 @@
|
||||||
|
SKIP: crbug.com/tint/938
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
Validation Failure:
|
||||||
|
|
||||||
|
Compilation failed:
|
||||||
|
|
||||||
|
program_source:5:24: error: no matching constructor for initialization of 'threadgroup metal::float2x3' (aka 'threadgroup matrix<float, 2, 3>')
|
||||||
|
threadgroup float2x3 tint_symbol_2;
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:269:23: note: candidate template ignored: requirement 'sizeof...(U) == 2' was not satisfied [with U = <>]
|
||||||
|
METAL_FUNC explicit matrix(initializer_list<U>... cols) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:281:23: note: candidate template ignored: requirement '(sizeof...(U) == 2) || (sizeof...(U) == 2 * 3)' was not satisfied [with U = <>]
|
||||||
|
METAL_FUNC explicit matrix(U... vals) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:299:23: note: candidate template ignored: requirement 'sizeof...(U) == 2' was not satisfied [with U = <>]
|
||||||
|
METAL_FUNC explicit matrix(initializer_list<U>... cols) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:311:23: note: candidate template ignored: requirement '(sizeof...(U) == 2) || (sizeof...(U) == 2 * 3)' was not satisfied [with U = <>]
|
||||||
|
METAL_FUNC explicit matrix(U... vals) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:330:23: note: candidate template ignored: requirement 'sizeof...(U) == 2' was not satisfied [with U = <>]
|
||||||
|
METAL_FUNC explicit matrix(initializer_list<U>... cols) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:342:23: note: candidate template ignored: requirement '(sizeof...(U) == 2) || (sizeof...(U) == 2 * 3)' was not satisfied [with U = <>]
|
||||||
|
METAL_FUNC explicit matrix(U... vals) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:56:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(T val, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:62:23: note: candidate constructor template not viable: requires 3 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(cols_init_tag, initializer_list<vec<T, Rows>> cols, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:68:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(cols_all_tag, U... cols) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:75:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(elems_all_tag, U... elems) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:80:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(initializer_list<T> elems, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:86:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(T val, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:92:23: note: candidate constructor template not viable: requires 3 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(cols_init_tag, initializer_list<vec<T, Rows>> cols, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:98:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(cols_all_tag, U... cols) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:105:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(elems_all_tag, U... elems) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:110:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(initializer_list<T> elems, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:117:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(T val, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:123:23: note: candidate constructor template not viable: requires 3 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(cols_init_tag, initializer_list<vec<T, Rows>> cols, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:129:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(cols_all_tag, U... cols) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:136:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(elems_all_tag, U... elems) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:141:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(initializer_list<T> elems, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:149:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:154:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:159:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:164:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:170:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:177:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:183:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:188:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:193:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:198:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:204:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:211:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:218:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:225:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:232:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:239:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:247:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:255:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
|
||||||
|
METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:358:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:365:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:372:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:379:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:387:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:396:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that) thread
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:404:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:411:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:418:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:425:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:433:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:442:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that) constant
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:451:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:460:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:469:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:478:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:488:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that) ray_data
|
||||||
|
^
|
||||||
|
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:498:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
|
||||||
|
METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that) ray_data
|
||||||
|
^
|
||||||
|
program_source:6:31: warning: equality comparison with extraneous parentheses
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
~~~~~~~~~~~~~~~~~~~~~~~^~~~~
|
||||||
|
program_source:6:31: note: remove extraneous parentheses around the comparison to silence this warning
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
~ ^ ~
|
||||||
|
program_source:6:31: note: use '=' to turn this equality comparison into an assignment
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
^~
|
||||||
|
=
|
||||||
|
|
@ -0,0 +1,42 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 25
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%v3float = OpTypeVector %float 3
|
||||||
|
%mat2v3float = OpTypeMatrix %v3float 2
|
||||||
|
%_ptr_Workgroup_mat2v3float = OpTypePointer Workgroup %mat2v3float
|
||||||
|
%v = OpVariable %_ptr_Workgroup_mat2v3float Workgroup
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_uint Input
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%9 = OpTypeFunction %void
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%19 = OpConstantNull %mat2v3float
|
||||||
|
%uint_2 = OpConstant %uint 2
|
||||||
|
%uint_264 = OpConstant %uint 264
|
||||||
|
%main = OpFunction %void None %9
|
||||||
|
%12 = OpLabel
|
||||||
|
%13 = OpLoad %uint %tint_symbol
|
||||||
|
%15 = OpIEqual %bool %13 %uint_0
|
||||||
|
OpSelectionMerge %17 None
|
||||||
|
OpBranchConditional %15 %18 %17
|
||||||
|
%18 = OpLabel
|
||||||
|
OpStore %v %19
|
||||||
|
OpBranch %17
|
||||||
|
%17 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
%24 = OpLoad %mat2v3float %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : mat2x3<f32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : i32;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,16 @@
|
||||||
|
groupshared int v;
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint local_invocation_index : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
v = 0;
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,13 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||||
|
threadgroup int tint_symbol_2;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
tint_symbol_2 = int();
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
(void) tint_symbol_2;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,40 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 23
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%_ptr_Workgroup_int = OpTypePointer Workgroup %int
|
||||||
|
%v = OpVariable %_ptr_Workgroup_int Workgroup
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_uint Input
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%7 = OpTypeFunction %void
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%17 = OpConstantNull %int
|
||||||
|
%uint_2 = OpConstant %uint 2
|
||||||
|
%uint_264 = OpConstant %uint 264
|
||||||
|
%main = OpFunction %void None %7
|
||||||
|
%10 = OpLabel
|
||||||
|
%11 = OpLoad %uint %tint_symbol
|
||||||
|
%13 = OpIEqual %bool %11 %uint_0
|
||||||
|
OpSelectionMerge %15 None
|
||||||
|
OpBranchConditional %13 %16 %15
|
||||||
|
%16 = OpLabel
|
||||||
|
OpStore %v %17
|
||||||
|
OpBranch %15
|
||||||
|
%15 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
%22 = OpLoad %int %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : i32;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,11 @@
|
||||||
|
struct S {
|
||||||
|
a : i32;
|
||||||
|
b : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
var<workgroup> v : S;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,22 @@
|
||||||
|
struct S {
|
||||||
|
int a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
groupshared S v;
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint local_invocation_index : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
const S tint_symbol_2 = {0, 0.0f};
|
||||||
|
v = tint_symbol_2;
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,19 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
struct S {
|
||||||
|
int a;
|
||||||
|
float b;
|
||||||
|
};
|
||||||
|
|
||||||
|
kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||||
|
threadgroup S tint_symbol_3;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
S const tint_symbol_2 = {};
|
||||||
|
tint_symbol_3 = tint_symbol_2;
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
(void) tint_symbol_3;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,47 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 25
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %S "S"
|
||||||
|
OpMemberName %S 0 "a"
|
||||||
|
OpMemberName %S 1 "b"
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpMemberDecorate %S 0 Offset 0
|
||||||
|
OpMemberDecorate %S 1 Offset 4
|
||||||
|
OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%float = OpTypeFloat 32
|
||||||
|
%S = OpTypeStruct %int %float
|
||||||
|
%_ptr_Workgroup_S = OpTypePointer Workgroup %S
|
||||||
|
%v = OpVariable %_ptr_Workgroup_S Workgroup
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_uint Input
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%9 = OpTypeFunction %void
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%19 = OpConstantNull %S
|
||||||
|
%uint_2 = OpConstant %uint 2
|
||||||
|
%uint_264 = OpConstant %uint 264
|
||||||
|
%main = OpFunction %void None %9
|
||||||
|
%12 = OpLabel
|
||||||
|
%13 = OpLoad %uint %tint_symbol
|
||||||
|
%15 = OpIEqual %bool %13 %uint_0
|
||||||
|
OpSelectionMerge %17 None
|
||||||
|
OpBranchConditional %15 %18 %17
|
||||||
|
%18 = OpLabel
|
||||||
|
OpStore %v %19
|
||||||
|
OpBranch %17
|
||||||
|
%17 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
%24 = OpLoad %S %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,11 @@
|
||||||
|
struct S {
|
||||||
|
a : i32;
|
||||||
|
b : f32;
|
||||||
|
};
|
||||||
|
|
||||||
|
var<workgroup> v : S;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : vec3<i32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,16 @@
|
||||||
|
groupshared int3 v;
|
||||||
|
|
||||||
|
struct tint_symbol_1 {
|
||||||
|
uint local_invocation_index : SV_GroupIndex;
|
||||||
|
};
|
||||||
|
|
||||||
|
[numthreads(1, 1, 1)]
|
||||||
|
void main(tint_symbol_1 tint_symbol) {
|
||||||
|
const uint local_invocation_index = tint_symbol.local_invocation_index;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
v = int3(0, 0, 0);
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
v;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
@ -0,0 +1,13 @@
|
||||||
|
#include <metal_stdlib>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
|
||||||
|
threadgroup int3 tint_symbol_2;
|
||||||
|
if ((local_invocation_index == 0u)) {
|
||||||
|
tint_symbol_2 = int3();
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
(void) tint_symbol_2;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
|
@ -0,0 +1,41 @@
|
||||||
|
; SPIR-V
|
||||||
|
; Version: 1.3
|
||||||
|
; Generator: Google Tint Compiler; 0
|
||||||
|
; Bound: 24
|
||||||
|
; Schema: 0
|
||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %v "v"
|
||||||
|
OpName %tint_symbol "tint_symbol"
|
||||||
|
OpName %main "main"
|
||||||
|
OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
|
||||||
|
%int = OpTypeInt 32 1
|
||||||
|
%v3int = OpTypeVector %int 3
|
||||||
|
%_ptr_Workgroup_v3int = OpTypePointer Workgroup %v3int
|
||||||
|
%v = OpVariable %_ptr_Workgroup_v3int Workgroup
|
||||||
|
%uint = OpTypeInt 32 0
|
||||||
|
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||||
|
%tint_symbol = OpVariable %_ptr_Input_uint Input
|
||||||
|
%void = OpTypeVoid
|
||||||
|
%8 = OpTypeFunction %void
|
||||||
|
%uint_0 = OpConstant %uint 0
|
||||||
|
%bool = OpTypeBool
|
||||||
|
%18 = OpConstantNull %v3int
|
||||||
|
%uint_2 = OpConstant %uint 2
|
||||||
|
%uint_264 = OpConstant %uint 264
|
||||||
|
%main = OpFunction %void None %8
|
||||||
|
%11 = OpLabel
|
||||||
|
%12 = OpLoad %uint %tint_symbol
|
||||||
|
%14 = OpIEqual %bool %12 %uint_0
|
||||||
|
OpSelectionMerge %16 None
|
||||||
|
OpBranchConditional %14 %17 %16
|
||||||
|
%17 = OpLabel
|
||||||
|
OpStore %v %18
|
||||||
|
OpBranch %16
|
||||||
|
%16 = OpLabel
|
||||||
|
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||||
|
%23 = OpLoad %v3int %v
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
|
|
@ -0,0 +1,6 @@
|
||||||
|
var<workgroup> v : vec3<i32>;
|
||||||
|
|
||||||
|
[[stage(compute), workgroup_size(1)]]
|
||||||
|
fn main() {
|
||||||
|
ignore(v);
|
||||||
|
}
|
||||||
Loading…
Reference in New Issue