test: Add bug cases taken from monorail
Some now pass, some still fail. If we fix bugs and they begin to pass, we have a better chance of noticing and marking the issues as fixed. Change-Id: I7278f960967a570ccde8865dcf2aa580235559bf Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54062 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
aba42ed362
commit
3628949efc
|
|
@ -1,35 +1,18 @@
|
|||
SKIP: crbug.com/tint/833 loop emission is broken
|
||||
SKIP: FAILED
|
||||
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Buf {
|
||||
/* 0x0000 */ uint count;
|
||||
/* 0x0004 */ uint data[50];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(device Buf& b [[buffer(0)]]) {
|
||||
uint i = 0u;
|
||||
{
|
||||
bool tint_msl_is_first_1 = true;
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:16:24: error: default initialization of an object of const type 'device uint *const' (aka 'device unsigned int *const')
|
||||
device uint* const p;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
*(p) = (*(p) * 2u);
|
||||
i = (i + 1u);
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
if ((i >= b.count)) {
|
||||
break;
|
||||
}
|
||||
p = &(b.data[i]);
|
||||
if (((i % 2u) == 0u)) {
|
||||
continue;
|
||||
}
|
||||
*(p) = 0u;
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
^
|
||||
= nullptr
|
||||
program_source:27:9: error: cannot assign to variable 'p' with const-qualified type 'device uint *const' (aka 'device unsigned int *const')
|
||||
p = &(b.data.array[i]);
|
||||
~ ^
|
||||
program_source:16:24: note: variable 'p' declared const here
|
||||
device uint* const p;
|
||||
~~~~~~~~~~~~~~~~~~~^
|
||||
|
|
|
|||
|
|
@ -0,0 +1,6 @@
|
|||
[[stage(vertex)]]
|
||||
fn main() {
|
||||
var light : vec3<f32> = vec3<f32>(1.2, 1., 2.);
|
||||
var negative_light : vec3<f32> = -light;
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,6 @@
|
|||
void main() {
|
||||
float3 light = float3(1.200000048f, 1.0f, 2.0f);
|
||||
float3 negative_light = -(light);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,9 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
vertex void tint_symbol() {
|
||||
float3 light = float3(1.200000048f, 1.0f, 2.0f);
|
||||
float3 negative_light = -(light);
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,37 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 20
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %main "main" %tint_pointsize
|
||||
OpName %tint_pointsize "tint_pointsize"
|
||||
OpName %main "main"
|
||||
OpName %light "light"
|
||||
OpName %negative_light "negative_light"
|
||||
OpDecorate %tint_pointsize BuiltIn PointSize
|
||||
%float = OpTypeFloat 32
|
||||
%_ptr_Output_float = OpTypePointer Output %float
|
||||
%4 = OpConstantNull %float
|
||||
%tint_pointsize = OpVariable %_ptr_Output_float Output %4
|
||||
%void = OpTypeVoid
|
||||
%5 = OpTypeFunction %void
|
||||
%float_1 = OpConstant %float 1
|
||||
%v3float = OpTypeVector %float 3
|
||||
%float_1_20000005 = OpConstant %float 1.20000005
|
||||
%float_2 = OpConstant %float 2
|
||||
%13 = OpConstantComposite %v3float %float_1_20000005 %float_1 %float_2
|
||||
%_ptr_Function_v3float = OpTypePointer Function %v3float
|
||||
%16 = OpConstantNull %v3float
|
||||
%main = OpFunction %void None %5
|
||||
%8 = OpLabel
|
||||
%light = OpVariable %_ptr_Function_v3float Function %16
|
||||
%negative_light = OpVariable %_ptr_Function_v3float Function %16
|
||||
OpStore %tint_pointsize %float_1
|
||||
OpStore %light %13
|
||||
%18 = OpLoad %v3float %light
|
||||
%17 = OpFNegate %v3float %18
|
||||
OpStore %negative_light %17
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,6 @@
|
|||
[[stage(vertex)]]
|
||||
fn main() {
|
||||
var light : vec3<f32> = vec3<f32>(1.200000048, 1.0, 2.0);
|
||||
var negative_light : vec3<f32> = -(light);
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,8 @@
|
|||
struct Light {
|
||||
position : vec3<f32>;
|
||||
colour : vec3<f32>;
|
||||
};
|
||||
[[block]] struct Lights {
|
||||
light : [[stride(32)]] array<Light>;
|
||||
};
|
||||
[[set(0), binding(1)]] var<storage> lights : [[access(read)]] Lights;
|
||||
|
|
@ -0,0 +1,10 @@
|
|||
bug/tint/294.wgsl:8:24 warning: use of deprecated language feature: declare access with var<storage, read> instead of using [[access]] decoration
|
||||
[[set(0), binding(1)]] var<storage> lights : [[access(read)]] Lights;
|
||||
^^^
|
||||
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,17 @@
|
|||
bug/tint/294.wgsl:8:24 warning: use of deprecated language feature: declare access with var<storage, read> instead of using [[access]] decoration
|
||||
[[set(0), binding(1)]] var<storage> lights : [[access(read)]] Lights;
|
||||
^^^
|
||||
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Light {
|
||||
/* 0x0000 */ packed_float3 position;
|
||||
/* 0x000c */ int8_t tint_pad_0[4];
|
||||
/* 0x0010 */ packed_float3 colour;
|
||||
/* 0x001c */ int8_t tint_pad_1[4];
|
||||
};
|
||||
struct Lights {
|
||||
/* 0x0000 */ Light light[1];
|
||||
};
|
||||
|
||||
|
|
@ -0,0 +1,41 @@
|
|||
bug/tint/294.wgsl:8:24 warning: use of deprecated language feature: declare access with var<storage, read> instead of using [[access]] decoration
|
||||
[[set(0), binding(1)]] var<storage> lights : [[access(read)]] Lights;
|
||||
^^^
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 12
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
|
||||
OpExecutionMode %unused_entry_point LocalSize 1 1 1
|
||||
OpName %Lights "Lights"
|
||||
OpMemberName %Lights 0 "light"
|
||||
OpName %Light "Light"
|
||||
OpMemberName %Light 0 "position"
|
||||
OpMemberName %Light 1 "colour"
|
||||
OpName %lights "lights"
|
||||
OpName %unused_entry_point "unused_entry_point"
|
||||
OpDecorate %Lights Block
|
||||
OpMemberDecorate %Lights 0 Offset 0
|
||||
OpMemberDecorate %Light 0 Offset 0
|
||||
OpMemberDecorate %Light 1 Offset 16
|
||||
OpDecorate %_runtimearr_Light ArrayStride 32
|
||||
OpDecorate %lights NonWritable
|
||||
OpDecorate %lights DescriptorSet 0
|
||||
OpDecorate %lights Binding 1
|
||||
%float = OpTypeFloat 32
|
||||
%v3float = OpTypeVector %float 3
|
||||
%Light = OpTypeStruct %v3float %v3float
|
||||
%_runtimearr_Light = OpTypeRuntimeArray %Light
|
||||
%Lights = OpTypeStruct %_runtimearr_Light
|
||||
%_ptr_StorageBuffer_Lights = OpTypePointer StorageBuffer %Lights
|
||||
%lights = OpVariable %_ptr_StorageBuffer_Lights StorageBuffer
|
||||
%void = OpTypeVoid
|
||||
%8 = OpTypeFunction %void
|
||||
%unused_entry_point = OpFunction %void None %8
|
||||
%11 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
bug/tint/294.wgsl:8:24 warning: use of deprecated language feature: declare access with var<storage, read> instead of using [[access]] decoration
|
||||
[[set(0), binding(1)]] var<storage> lights : [[access(read)]] Lights;
|
||||
^^^
|
||||
|
||||
struct Light {
|
||||
position : vec3<f32>;
|
||||
colour : vec3<f32>;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct Lights {
|
||||
light : [[stride(32)]] array<Light>;
|
||||
};
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> lights : Lights;
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
|
||||
[[block]]
|
||||
struct vertexUniformBuffer1 {
|
||||
transform1 : mat2x2<f32>;
|
||||
};
|
||||
[[block]]
|
||||
struct vertexUniformBuffer2 {
|
||||
transform2 : mat2x2<f32>;
|
||||
};
|
||||
|
||||
[[set(0), binding(0)]] var<uniform> x_20 : vertexUniformBuffer1;
|
||||
[[set(1), binding(0)]] var<uniform> x_26 : vertexUniformBuffer2;
|
||||
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn main([[builtin(vertex_idx)]] gl_VertexIndex : i32) -> [[builtin(position)]] vec4<f32> {
|
||||
var indexable : array<vec2<f32>, 3>;
|
||||
let x_23 : mat2x2<f32> = x_20.transform1;
|
||||
let x_28 : mat2x2<f32> = x_26.transform2;
|
||||
let x_46 : i32 = gl_VertexIndex;
|
||||
indexable = array<vec2<f32>, 3>(vec2<f32>(-1.0, 1.0), vec2<f32>(1.0, 1.0), vec2<f32>(-1.0, -1.0));
|
||||
let x_51 : vec2<f32> = indexable[x_46];
|
||||
let x_52 : vec2<f32> = (mat2x2<f32>((x_23[0u] + x_28[0u]), (x_23[1u] + x_28[1u])) * x_51);
|
||||
return vec4<f32>(x_52.x, x_52.y, 0.0, 1.0);
|
||||
}
|
||||
|
|
@ -0,0 +1,43 @@
|
|||
SKIP: FAILED
|
||||
|
||||
|
||||
|
||||
Validation Failure:
|
||||
struct vertexUniformBuffer1 {
|
||||
float2x2 transform1;
|
||||
};
|
||||
struct vertexUniformBuffer2 {
|
||||
float2x2 transform2;
|
||||
};
|
||||
struct tint_symbol_1 {
|
||||
int gl_VertexIndex : SV_VertexID;
|
||||
};
|
||||
struct tint_symbol_2 {
|
||||
float4 value : SV_Position;
|
||||
};
|
||||
|
||||
ConstantBuffer<vertexUniformBuffer1> x_20 : register(b0, space0);
|
||||
ConstantBuffer<vertexUniformBuffer2> x_26 : register(b0, space1);
|
||||
|
||||
tint_symbol_2 main(tint_symbol_1 tint_symbol) {
|
||||
const int gl_VertexIndex = tint_symbol.gl_VertexIndex;
|
||||
float2 indexable[3] = {float2(0.0f, 0.0f), float2(0.0f, 0.0f), float2(0.0f, 0.0f)};
|
||||
const float2x2 x_23 = x_20.transform1;
|
||||
const float2x2 x_28 = x_26.transform2;
|
||||
const int x_46 = gl_VertexIndex;
|
||||
const float2 tint_symbol_3[3] = {float2(-1.0f, 1.0f), float2(1.0f, 1.0f), float2(-1.0f, -1.0f)};
|
||||
indexable = tint_symbol_3;
|
||||
const float2 x_51 = indexable[x_46];
|
||||
const float2 x_52 = mul(x_51, float2x2((x_23[0u] + x_28[0u]), (x_23[1u] + x_28[1u])));
|
||||
const tint_symbol_2 tint_symbol_4 = {float4(x_52.x, x_52.y, 0.0f, 1.0f)};
|
||||
return tint_symbol_4;
|
||||
}
|
||||
|
||||
|
||||
warning: DXIL.dll not found. Resulting DXIL will not be signed for use in release environments.
|
||||
|
||||
error: validation errors
|
||||
error: SV_VertexID must be uint.
|
||||
Validation failed.
|
||||
|
||||
|
||||
|
|
@ -0,0 +1,19 @@
|
|||
SKIP: FAILED
|
||||
|
||||
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:17:55: error: type 'int' is not valid for attribute 'vertex_id'
|
||||
vertex tint_symbol_2 tint_symbol(int gl_VertexIndex [[vertex_id]], constant vertexUniformBuffer1& x_20 [[buffer(0)]], constant vertexUniformBuffer2& x_26 [[buffer(0)]]) {
|
||||
^~~~~~~~~
|
||||
program_source:18:37: warning: suggest braces around initialization of subobject
|
||||
tint_array_wrapper_0 indexable = {0.0f};
|
||||
^~~~
|
||||
{ }
|
||||
program_source:22:47: warning: suggest braces around initialization of subobject
|
||||
tint_array_wrapper_0 const tint_symbol_3 = {float2(-1.0f, 1.0f), float2(1.0f, 1.0f), float2(-1.0f, -1.0f)};
|
||||
^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
|
||||
{ }
|
||||
|
|
@ -0,0 +1,109 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 64
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Vertex %main "main" %tint_pointsize %tint_symbol %tint_symbol_2
|
||||
OpName %tint_pointsize "tint_pointsize"
|
||||
OpName %vertexUniformBuffer1 "vertexUniformBuffer1"
|
||||
OpMemberName %vertexUniformBuffer1 0 "transform1"
|
||||
OpName %x_20 "x_20"
|
||||
OpName %vertexUniformBuffer2 "vertexUniformBuffer2"
|
||||
OpMemberName %vertexUniformBuffer2 0 "transform2"
|
||||
OpName %x_26 "x_26"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %tint_symbol_2 "tint_symbol_2"
|
||||
OpName %tint_symbol_3 "tint_symbol_3"
|
||||
OpName %tint_symbol_1 "tint_symbol_1"
|
||||
OpName %main "main"
|
||||
OpName %indexable "indexable"
|
||||
OpDecorate %tint_pointsize BuiltIn PointSize
|
||||
OpDecorate %vertexUniformBuffer1 Block
|
||||
OpMemberDecorate %vertexUniformBuffer1 0 Offset 0
|
||||
OpMemberDecorate %vertexUniformBuffer1 0 ColMajor
|
||||
OpMemberDecorate %vertexUniformBuffer1 0 MatrixStride 8
|
||||
OpDecorate %x_20 NonWritable
|
||||
OpDecorate %x_20 DescriptorSet 0
|
||||
OpDecorate %x_20 Binding 0
|
||||
OpDecorate %vertexUniformBuffer2 Block
|
||||
OpMemberDecorate %vertexUniformBuffer2 0 Offset 0
|
||||
OpMemberDecorate %vertexUniformBuffer2 0 ColMajor
|
||||
OpMemberDecorate %vertexUniformBuffer2 0 MatrixStride 8
|
||||
OpDecorate %x_26 NonWritable
|
||||
OpDecorate %x_26 DescriptorSet 1
|
||||
OpDecorate %x_26 Binding 0
|
||||
OpDecorate %tint_symbol BuiltIn VertexIndex
|
||||
OpDecorate %tint_symbol_2 BuiltIn Position
|
||||
OpDecorate %_arr_v2float_uint_3 ArrayStride 8
|
||||
%float = OpTypeFloat 32
|
||||
%_ptr_Output_float = OpTypePointer Output %float
|
||||
%4 = OpConstantNull %float
|
||||
%tint_pointsize = OpVariable %_ptr_Output_float Output %4
|
||||
%v2float = OpTypeVector %float 2
|
||||
%mat2v2float = OpTypeMatrix %v2float 2
|
||||
%vertexUniformBuffer1 = OpTypeStruct %mat2v2float
|
||||
%_ptr_Uniform_vertexUniformBuffer1 = OpTypePointer Uniform %vertexUniformBuffer1
|
||||
%x_20 = OpVariable %_ptr_Uniform_vertexUniformBuffer1 Uniform
|
||||
%vertexUniformBuffer2 = OpTypeStruct %mat2v2float
|
||||
%_ptr_Uniform_vertexUniformBuffer2 = OpTypePointer Uniform %vertexUniformBuffer2
|
||||
%x_26 = OpVariable %_ptr_Uniform_vertexUniformBuffer2 Uniform
|
||||
%int = OpTypeInt 32 1
|
||||
%_ptr_Input_int = OpTypePointer Input %int
|
||||
%tint_symbol = OpVariable %_ptr_Input_int Input
|
||||
%v4float = OpTypeVector %float 4
|
||||
%_ptr_Output_v4float = OpTypePointer Output %v4float
|
||||
%19 = OpConstantNull %v4float
|
||||
%tint_symbol_2 = OpVariable %_ptr_Output_v4float Output %19
|
||||
%void = OpTypeVoid
|
||||
%20 = OpTypeFunction %void %v4float
|
||||
%25 = OpTypeFunction %void
|
||||
%float_1 = OpConstant %float 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_arr_v2float_uint_3 = OpTypeArray %v2float %uint_3
|
||||
%_ptr_Function__arr_v2float_uint_3 = OpTypePointer Function %_arr_v2float_uint_3
|
||||
%34 = OpConstantNull %_arr_v2float_uint_3
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
|
||||
%float_n1 = OpConstant %float -1
|
||||
%43 = OpConstantComposite %v2float %float_n1 %float_1
|
||||
%44 = OpConstantComposite %v2float %float_1 %float_1
|
||||
%45 = OpConstantComposite %v2float %float_n1 %float_n1
|
||||
%46 = OpConstantComposite %_arr_v2float_uint_3 %43 %44 %45
|
||||
%_ptr_Function_v2float = OpTypePointer Function %v2float
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%float_0 = OpConstant %float 0
|
||||
%tint_symbol_3 = OpFunction %void None %20
|
||||
%tint_symbol_1 = OpFunctionParameter %v4float
|
||||
%24 = OpLabel
|
||||
OpStore %tint_symbol_2 %tint_symbol_1
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %25
|
||||
%27 = OpLabel
|
||||
%indexable = OpVariable %_ptr_Function__arr_v2float_uint_3 Function %34
|
||||
OpStore %tint_pointsize %float_1
|
||||
%37 = OpAccessChain %_ptr_Uniform_mat2v2float %x_20 %uint_0
|
||||
%38 = OpLoad %mat2v2float %37
|
||||
%39 = OpAccessChain %_ptr_Uniform_mat2v2float %x_26 %uint_0
|
||||
%40 = OpLoad %mat2v2float %39
|
||||
%41 = OpLoad %int %tint_symbol
|
||||
OpStore %indexable %46
|
||||
%48 = OpAccessChain %_ptr_Function_v2float %indexable %41
|
||||
%49 = OpLoad %v2float %48
|
||||
%50 = OpCompositeExtract %v2float %38 0
|
||||
%51 = OpCompositeExtract %v2float %40 0
|
||||
%52 = OpFAdd %v2float %50 %51
|
||||
%54 = OpCompositeExtract %v2float %38 1
|
||||
%55 = OpCompositeExtract %v2float %40 1
|
||||
%56 = OpFAdd %v2float %54 %55
|
||||
%57 = OpCompositeConstruct %mat2v2float %52 %56
|
||||
%58 = OpMatrixTimesVector %v2float %57 %49
|
||||
%60 = OpCompositeExtract %float %58 0
|
||||
%61 = OpCompositeExtract %float %58 1
|
||||
%63 = OpCompositeConstruct %v4float %60 %61 %float_0 %float_1
|
||||
%59 = OpFunctionCall %void %tint_symbol_3 %63
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
[[block]]
|
||||
struct vertexUniformBuffer1 {
|
||||
transform1 : mat2x2<f32>;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct vertexUniformBuffer2 {
|
||||
transform2 : mat2x2<f32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> x_20 : vertexUniformBuffer1;
|
||||
|
||||
[[group(1), binding(0)]] var<uniform> x_26 : vertexUniformBuffer2;
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn main([[builtin(vertex_index)]] gl_VertexIndex : i32) -> [[builtin(position)]] vec4<f32> {
|
||||
var indexable : array<vec2<f32>, 3>;
|
||||
let x_23 : mat2x2<f32> = x_20.transform1;
|
||||
let x_28 : mat2x2<f32> = x_26.transform2;
|
||||
let x_46 : i32 = gl_VertexIndex;
|
||||
indexable = array<vec2<f32>, 3>(vec2<f32>(-1.0, 1.0), vec2<f32>(1.0, 1.0), vec2<f32>(-1.0, -1.0));
|
||||
let x_51 : vec2<f32> = indexable[x_46];
|
||||
let x_52 : vec2<f32> = (mat2x2<f32>((x_23[0u] + x_28[0u]), (x_23[1u] + x_28[1u])) * x_51);
|
||||
return vec4<f32>(x_52.x, x_52.y, 0.0, 1.0);
|
||||
}
|
||||
|
|
@ -0,0 +1,50 @@
|
|||
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource GLSL 450
|
||||
OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
|
||||
OpSourceExtension "GL_GOOGLE_include_directive"
|
||||
OpName %main "main"
|
||||
OpName %srcValue "srcValue"
|
||||
OpName %Src "Src"
|
||||
OpName %Dst "Dst"
|
||||
OpDecorate %Src DescriptorSet 0
|
||||
OpDecorate %Src Binding 0
|
||||
OpDecorate %Src NonWritable
|
||||
OpDecorate %Dst DescriptorSet 0
|
||||
OpDecorate %Dst Binding 1
|
||||
OpDecorate %Dst NonReadable
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%v4uint = OpTypeVector %uint 4
|
||||
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
|
||||
%10 = OpTypeImage %uint 2D 0 0 0 2 R32ui
|
||||
%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
|
||||
%Src = OpVariable %_ptr_UniformConstant_10 UniformConstant
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%17 = OpConstantComposite %v2int %int_0 %int_0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%int_1 = OpConstant %int 1
|
||||
%Dst = OpVariable %_ptr_UniformConstant_10 UniformConstant
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%srcValue = OpVariable %_ptr_Function_v4uint Function
|
||||
%13 = OpLoad %10 %Src
|
||||
%18 = OpImageRead %v4uint %13 %17
|
||||
OpStore %srcValue %18
|
||||
%21 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
|
||||
%22 = OpLoad %uint %21
|
||||
%24 = OpIAdd %uint %22 %int_1
|
||||
OpStore %21 %24
|
||||
%26 = OpLoad %10 %Dst
|
||||
%27 = OpLoad %v4uint %srcValue
|
||||
OpImageWrite %26 %17 %27
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
Texture2D<uint4> Src : register(t0, space0);
|
||||
RWTexture2D<uint4> Dst : register(u1, space0);
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main() {
|
||||
uint4 srcValue = uint4(0u, 0u, 0u, 0u);
|
||||
const uint4 x_18 = Src.Load(int3(0, 0, 0));
|
||||
srcValue = x_18;
|
||||
const uint x_22 = srcValue.x;
|
||||
srcValue.x = (x_22 + asuint(1));
|
||||
const uint4 x_27 = srcValue;
|
||||
Dst[int2(0, 0)] = x_27;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,17 @@
|
|||
SKIP: FAILED
|
||||
|
||||
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:6:22: error: use of undeclared identifier 'Src'
|
||||
uint4 const x_18 = Src.read(int2(0, 0));
|
||||
^
|
||||
program_source:8:29: error: address of vector element requested
|
||||
thread uint* const x_21 = &(srcValue.x);
|
||||
^ ~~~~~~~~~~
|
||||
program_source:12:3: error: use of undeclared identifier 'Dst'
|
||||
Dst.write(x_27, int2(0, 0));
|
||||
^
|
||||
|
|
@ -0,0 +1,56 @@
|
|||
SKIP: FAILED
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 32
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main"
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Src "Src"
|
||||
OpName %Dst "Dst"
|
||||
OpName %main "main"
|
||||
OpName %srcValue "srcValue"
|
||||
OpDecorate %Src NonWritable
|
||||
OpDecorate %Src DescriptorSet 0
|
||||
OpDecorate %Src Binding 0
|
||||
OpDecorate %Dst NonReadable
|
||||
OpDecorate %Dst DescriptorSet 0
|
||||
OpDecorate %Dst Binding 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%3 = OpTypeImage %uint 2D 0 0 0 2 R32ui
|
||||
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
|
||||
%Src = OpVariable %_ptr_UniformConstant_3 UniformConstant
|
||||
%_ptr_UniformConstant_3_0 = OpTypePointer UniformConstant %3
|
||||
%Dst = OpVariable %_ptr_UniformConstant_3_0 UniformConstant
|
||||
%void = OpTypeVoid
|
||||
%7 = OpTypeFunction %void
|
||||
%v4uint = OpTypeVector %uint 4
|
||||
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
|
||||
%14 = OpConstantNull %v4uint
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%20 = OpConstantComposite %v2int %int_0 %int_0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%int_1 = OpConstant %int 1
|
||||
%main = OpFunction %void None %7
|
||||
%10 = OpLabel
|
||||
%srcValue = OpVariable %_ptr_Function_v4uint Function %14
|
||||
%16 = OpLoad %3 %Src
|
||||
%15 = OpImageRead %v4uint %16 %20
|
||||
OpStore %srcValue %15
|
||||
%23 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
|
||||
%24 = OpLoad %uint %23
|
||||
%25 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
|
||||
%26 = OpBitcast %uint %int_1
|
||||
%28 = OpIAdd %uint %24 %26
|
||||
OpStore %25 %28
|
||||
%29 = OpLoad %v4uint %srcValue
|
||||
%31 = OpLoad %3 %Dst
|
||||
OpImageWrite %31 %20 %29
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,16 @@
|
|||
[[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>;
|
||||
|
||||
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main() {
|
||||
var srcValue : vec4<u32>;
|
||||
let x_18 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
|
||||
srcValue = x_18;
|
||||
let x_21 : ptr<function, u32> = &(srcValue.x);
|
||||
let x_22 : u32 = *(x_21);
|
||||
*(x_21) = (x_22 + bitcast<u32>(1));
|
||||
let x_27 : vec4<u32> = srcValue;
|
||||
textureStore(Dst, vec2<i32>(0, 0), x_27);
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
|
||||
[[group(0), binding(1)]] var Dst : [[access(write)]] texture_storage_2d<r32uint>;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main() {
|
||||
var srcValue : vec4<u32>;
|
||||
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
|
||||
srcValue = x_22;
|
||||
let x_23 : ptr<function, u32> = &srcValue.x;
|
||||
let x_24 : u32 = *x_23;
|
||||
let x_25 : u32 = (x_24 + 1u);
|
||||
let x_27 : vec4<u32> = srcValue;
|
||||
textureStore(Dst, vec2<i32>(0, 0), x_27.xxxx);
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,23 @@
|
|||
bug/tint/453.wgsl:1:79 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
bug/tint/453.wgsl:2:80 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(1)]] var Dst : [[access(write)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
Texture2D<uint4> Src : register(t0, space0);
|
||||
RWTexture2D<uint4> Dst : register(u1, space0);
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main() {
|
||||
uint4 srcValue = uint4(0u, 0u, 0u, 0u);
|
||||
const uint4 x_22 = Src.Load(int3(0, 0, 0));
|
||||
srcValue = x_22;
|
||||
const uint x_24 = srcValue.x;
|
||||
const uint x_25 = (x_24 + 1u);
|
||||
const uint4 x_27 = srcValue;
|
||||
Dst[int2(0, 0)] = x_27.xxxx;
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,25 @@
|
|||
SKIP: FAILED
|
||||
|
||||
bug/tint/453.wgsl:1:79 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
bug/tint/453.wgsl:2:80 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(1)]] var Dst : [[access(write)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:6:22: error: use of undeclared identifier 'Src'
|
||||
uint4 const x_22 = Src.read(int2(0, 0));
|
||||
^
|
||||
program_source:8:29: error: address of vector element requested
|
||||
thread uint* const x_23 = &(srcValue.x);
|
||||
^ ~~~~~~~~~~
|
||||
program_source:12:3: error: use of undeclared identifier 'Dst'
|
||||
Dst.write(x_27.xxxx, int2(0, 0));
|
||||
^
|
||||
|
|
@ -0,0 +1,62 @@
|
|||
SKIP: FAILED
|
||||
|
||||
bug/tint/453.wgsl:1:79 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
bug/tint/453.wgsl:2:80 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(1)]] var Dst : [[access(write)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
; 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 %Src "Src"
|
||||
OpName %Dst "Dst"
|
||||
OpName %main "main"
|
||||
OpName %srcValue "srcValue"
|
||||
OpDecorate %Src NonWritable
|
||||
OpDecorate %Src DescriptorSet 0
|
||||
OpDecorate %Src Binding 0
|
||||
OpDecorate %Dst NonReadable
|
||||
OpDecorate %Dst DescriptorSet 0
|
||||
OpDecorate %Dst Binding 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%3 = OpTypeImage %uint 2D 0 0 0 2 R32ui
|
||||
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
|
||||
%Src = OpVariable %_ptr_UniformConstant_3 UniformConstant
|
||||
%_ptr_UniformConstant_3_0 = OpTypePointer UniformConstant %3
|
||||
%Dst = OpVariable %_ptr_UniformConstant_3_0 UniformConstant
|
||||
%void = OpTypeVoid
|
||||
%7 = OpTypeFunction %void
|
||||
%v4uint = OpTypeVector %uint 4
|
||||
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
|
||||
%14 = OpConstantNull %v4uint
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%20 = OpConstantComposite %v2int %int_0 %int_0
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%main = OpFunction %void None %7
|
||||
%10 = OpLabel
|
||||
%srcValue = OpVariable %_ptr_Function_v4uint Function %14
|
||||
%16 = OpLoad %3 %Src
|
||||
%15 = OpImageRead %v4uint %16 %20
|
||||
OpStore %srcValue %15
|
||||
%23 = OpAccessChain %_ptr_Function_uint %srcValue %uint_0
|
||||
%24 = OpLoad %uint %23
|
||||
%26 = OpIAdd %uint %24 %uint_1
|
||||
%27 = OpLoad %v4uint %srcValue
|
||||
%29 = OpLoad %3 %Dst
|
||||
%30 = OpVectorShuffle %v4uint %27 %27 0 0 0 0
|
||||
OpImageWrite %29 %20 %30
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,24 @@
|
|||
bug/tint/453.wgsl:1:79 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(0)]] var Src : [[access(read)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
bug/tint/453.wgsl:2:80 warning: use of deprecated language feature: access control is expected as last parameter of storage textures
|
||||
[[group(0), binding(1)]] var Dst : [[access(write)]] texture_storage_2d<r32uint>;
|
||||
^
|
||||
|
||||
[[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>;
|
||||
|
||||
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main() {
|
||||
var srcValue : vec4<u32>;
|
||||
let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
|
||||
srcValue = x_22;
|
||||
let x_23 : ptr<function, u32> = &(srcValue.x);
|
||||
let x_24 : u32 = *(x_23);
|
||||
let x_25 : u32 = (x_24 + 1u);
|
||||
let x_27 : vec4<u32> = srcValue;
|
||||
textureStore(Dst, vec2<i32>(0, 0), x_27.xxxx);
|
||||
return;
|
||||
}
|
||||
|
|
@ -1,49 +1,28 @@
|
|||
SKIP: crbug.com/tint/833 loop emission is broken
|
||||
SKIP: FAILED
|
||||
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ packed_uint2 aShape;
|
||||
/* 0x0008 */ packed_uint2 bShape;
|
||||
/* 0x0010 */ packed_uint2 outShape;
|
||||
};
|
||||
struct Matrix {
|
||||
/* 0x0000 */ uint numbers[1];
|
||||
};
|
||||
struct tint_symbol_2 {
|
||||
uint3 global_id [[thread_position_in_grid]];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
|
||||
uint3 const global_id = tint_symbol_1.global_id;
|
||||
uint2 const resultCell = uint2(global_id.y, global_id.x);
|
||||
uint const dimInner = uniforms.aShape.y;
|
||||
uint const dimOutter = uniforms.outShape.y;
|
||||
uint result = 0u;
|
||||
{
|
||||
uint i = 0u;
|
||||
{
|
||||
bool tint_msl_is_first_1 = true;
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:22:18: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const a;
|
||||
^
|
||||
= 0
|
||||
program_source:23:18: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const b;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
i = (i + 1u);
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
if (!((i < dimInner))) {
|
||||
break;
|
||||
}
|
||||
^
|
||||
= 0
|
||||
program_source:33:11: error: cannot assign to variable 'a' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
a = (i + (resultCell.x * dimInner));
|
||||
~ ^
|
||||
program_source:22:18: note: variable 'a' declared const here
|
||||
uint const a;
|
||||
~~~~~~~~~~~^
|
||||
program_source:34:11: error: cannot assign to variable 'b' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
b = (resultCell.y + (i * dimOutter));
|
||||
result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
|
||||
}
|
||||
}
|
||||
}
|
||||
uint const index = (resultCell.y + (resultCell.x * dimOutter));
|
||||
resultMatrix.numbers[index] = result;
|
||||
return;
|
||||
}
|
||||
|
||||
~ ^
|
||||
program_source:23:18: note: variable 'b' declared const here
|
||||
uint const b;
|
||||
~~~~~~~~~~~^
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -0,0 +1,26 @@
|
|||
|
||||
[[block]] struct Constants {
|
||||
[[offset(0)]] level : i32;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> constants : Constants;
|
||||
[[group(0), binding(1)]] var myTexture : texture_2d_array<f32>;
|
||||
|
||||
[[block]] struct Result {
|
||||
[[offset(0)]] values : [[stride(4)]] array<f32>;
|
||||
};
|
||||
[[group(0), binding(3)]] var<storage, read_write> result : Result;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
var flatIndex : u32 =
|
||||
2u * 2u * GlobalInvocationID.z +
|
||||
2u * GlobalInvocationID.y +
|
||||
GlobalInvocationID.x;
|
||||
flatIndex = flatIndex * 1u;
|
||||
var texel : vec4<f32> = textureLoad(myTexture, vec2<i32>(GlobalInvocationID.xy), 0, 0);
|
||||
|
||||
for (var i : u32 = 0u; i < 1u; i = i + 1u) {
|
||||
result.values[flatIndex + i] = texel.r;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,40 @@
|
|||
bug/tint/757.wgsl:3:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] level : i32;
|
||||
^^^^^^
|
||||
|
||||
bug/tint/757.wgsl:10:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] values : [[stride(4)]] array<f32>;
|
||||
^^^^^^
|
||||
|
||||
struct Constants {
|
||||
int level;
|
||||
};
|
||||
struct tint_symbol_1 {
|
||||
uint3 GlobalInvocationID : SV_DispatchThreadID;
|
||||
};
|
||||
|
||||
RWByteAddressBuffer result : register(u3, space0);
|
||||
|
||||
Texture2DArray<float4> myTexture : register(t1, space0);
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
const uint3 GlobalInvocationID = tint_symbol.GlobalInvocationID;
|
||||
uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
|
||||
flatIndex = (flatIndex * 1u);
|
||||
float4 texel = myTexture.Load(int4(GlobalInvocationID.xy, 0, 0), 0);
|
||||
{
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
if (!((i < 1u))) {
|
||||
break;
|
||||
}
|
||||
result.Store((4u * (flatIndex + i)), asuint(texel.r));
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,19 @@
|
|||
SKIP: FAILED
|
||||
|
||||
bug/tint/757.wgsl:3:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] level : i32;
|
||||
^^^^^^
|
||||
|
||||
bug/tint/757.wgsl:10:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] values : [[stride(4)]] array<f32>;
|
||||
^^^^^^
|
||||
|
||||
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:14:18: error: use of undeclared identifier 'myTexture'
|
||||
float4 texel = myTexture.read(int2(GlobalInvocationID.xy), 0, 0);
|
||||
^
|
||||
|
|
@ -0,0 +1,141 @@
|
|||
SKIP: FAILED
|
||||
|
||||
bug/tint/757.wgsl:3:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] level : i32;
|
||||
^^^^^^
|
||||
|
||||
bug/tint/757.wgsl:10:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] values : [[stride(4)]] array<f32>;
|
||||
^^^^^^
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 76
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %Constants "Constants"
|
||||
OpMemberName %Constants 0 "level"
|
||||
OpName %constants "constants"
|
||||
OpName %myTexture "myTexture"
|
||||
OpName %Result "Result"
|
||||
OpMemberName %Result 0 "values"
|
||||
OpName %result "result"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %main "main"
|
||||
OpName %flatIndex "flatIndex"
|
||||
OpName %texel "texel"
|
||||
OpName %i "i"
|
||||
OpDecorate %Constants Block
|
||||
OpMemberDecorate %Constants 0 Offset 0
|
||||
OpDecorate %constants NonWritable
|
||||
OpDecorate %constants DescriptorSet 0
|
||||
OpDecorate %constants Binding 0
|
||||
OpDecorate %myTexture DescriptorSet 0
|
||||
OpDecorate %myTexture Binding 1
|
||||
OpDecorate %Result Block
|
||||
OpMemberDecorate %Result 0 Offset 0
|
||||
OpDecorate %_runtimearr_float ArrayStride 4
|
||||
OpDecorate %result DescriptorSet 0
|
||||
OpDecorate %result Binding 3
|
||||
OpDecorate %tint_symbol BuiltIn GlobalInvocationId
|
||||
%int = OpTypeInt 32 1
|
||||
%Constants = OpTypeStruct %int
|
||||
%_ptr_Uniform_Constants = OpTypePointer Uniform %Constants
|
||||
%constants = OpVariable %_ptr_Uniform_Constants Uniform
|
||||
%float = OpTypeFloat 32
|
||||
%7 = OpTypeImage %float 2D 0 1 0 1 Unknown
|
||||
%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
|
||||
%myTexture = OpVariable %_ptr_UniformConstant_7 UniformConstant
|
||||
%_runtimearr_float = OpTypeRuntimeArray %float
|
||||
%Result = OpTypeStruct %_runtimearr_float
|
||||
%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
|
||||
%result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||
%void = OpTypeVoid
|
||||
%17 = OpTypeFunction %void
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%38 = OpConstantNull %uint
|
||||
%v4float = OpTypeVector %float 4
|
||||
%v3int = OpTypeVector %int 3
|
||||
%v2uint = OpTypeVector %uint 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Function_v4float = OpTypePointer Function %v4float
|
||||
%54 = OpConstantNull %v4float
|
||||
%bool = OpTypeBool
|
||||
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%main = OpFunction %void None %17
|
||||
%20 = OpLabel
|
||||
%flatIndex = OpVariable %_ptr_Function_uint Function %38
|
||||
%texel = OpVariable %_ptr_Function_v4float Function %54
|
||||
%i = OpVariable %_ptr_Function_uint Function %38
|
||||
%22 = OpIMul %uint %uint_2 %uint_2
|
||||
%24 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_2
|
||||
%25 = OpLoad %uint %24
|
||||
%26 = OpIMul %uint %22 %25
|
||||
%28 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1
|
||||
%29 = OpLoad %uint %28
|
||||
%30 = OpIMul %uint %uint_2 %29
|
||||
%31 = OpIAdd %uint %26 %30
|
||||
%33 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0
|
||||
%34 = OpLoad %uint %33
|
||||
%35 = OpIAdd %uint %31 %34
|
||||
OpStore %flatIndex %35
|
||||
%39 = OpLoad %uint %flatIndex
|
||||
%40 = OpIMul %uint %39 %uint_1
|
||||
OpStore %flatIndex %40
|
||||
%43 = OpLoad %7 %myTexture
|
||||
%46 = OpLoad %v3uint %tint_symbol
|
||||
%47 = OpVectorShuffle %v2uint %46 %46 0 1
|
||||
%48 = OpCompositeExtract %uint %47 0
|
||||
%49 = OpCompositeExtract %uint %47 1
|
||||
%51 = OpCompositeConstruct %v3int %48 %49 %int_0
|
||||
%41 = OpImageFetch %v4float %43 %51 Lod %int_0
|
||||
OpStore %texel %41
|
||||
OpStore %i %uint_0
|
||||
OpBranch %56
|
||||
%56 = OpLabel
|
||||
OpLoopMerge %57 %58 None
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
%61 = OpLoad %uint %i
|
||||
%62 = OpULessThan %bool %61 %uint_1
|
||||
%60 = OpLogicalNot %bool %62
|
||||
OpSelectionMerge %64 None
|
||||
OpBranchConditional %60 %65 %64
|
||||
%65 = OpLabel
|
||||
OpBranch %57
|
||||
%64 = OpLabel
|
||||
%66 = OpLoad %uint %flatIndex
|
||||
%67 = OpLoad %uint %i
|
||||
%68 = OpIAdd %uint %66 %67
|
||||
%70 = OpAccessChain %_ptr_StorageBuffer_float %result %uint_0 %68
|
||||
%72 = OpAccessChain %_ptr_Function_float %texel %uint_0
|
||||
%73 = OpLoad %float %72
|
||||
OpStore %70 %73
|
||||
OpBranch %58
|
||||
%58 = OpLabel
|
||||
%74 = OpLoad %uint %i
|
||||
%75 = OpIAdd %uint %74 %uint_1
|
||||
OpStore %i %75
|
||||
OpBranch %56
|
||||
%57 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
||||
|
||||
Validation Failure:
|
||||
1:1: Expected Constituents to be scalars or vectors of the same type as Result Type components
|
||||
%51 = OpCompositeConstruct %v3int %48 %49 %int_0
|
||||
|
||||
|
|
@ -0,0 +1,43 @@
|
|||
bug/tint/757.wgsl:3:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] level : i32;
|
||||
^^^^^^
|
||||
|
||||
bug/tint/757.wgsl:10:5 warning: use of deprecated language feature: [[offset]] has been replaced with [[size]] and [[align]]
|
||||
[[offset(0)]] values : [[stride(4)]] array<f32>;
|
||||
^^^^^^
|
||||
|
||||
[[block]]
|
||||
struct Constants {
|
||||
level : i32;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<uniform> constants : Constants;
|
||||
|
||||
[[group(0), binding(1)]] var myTexture : texture_2d_array<f32>;
|
||||
|
||||
[[block]]
|
||||
struct Result {
|
||||
values : [[stride(4)]] array<f32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(3)]] var<storage, read_write> result : Result;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
var flatIndex : u32 = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
|
||||
flatIndex = (flatIndex * 1u);
|
||||
var texel : vec4<f32> = textureLoad(myTexture, vec2<i32>(GlobalInvocationID.xy), 0, 0);
|
||||
{
|
||||
var i : u32 = 0u;
|
||||
loop {
|
||||
if (!((i < 1u))) {
|
||||
break;
|
||||
}
|
||||
result.values[(flatIndex + i)] = texel.r;
|
||||
|
||||
continuing {
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,10 @@
|
|||
fn f() {
|
||||
let m: mat4x4<f32> = mat4x4<f32>(
|
||||
vec4<f32>(1.0, 1.0, 1.0, 1.0),
|
||||
vec4<f32>(1.0, 1.0, 1.0, 1.0),
|
||||
vec4<f32>(1.0, 1.0, 1.0, 1.0),
|
||||
vec4<f32>(1.0, 1.0, 1.0, 1.0)
|
||||
);
|
||||
let v1: vec4<f32> = m[0];
|
||||
let a: f32 = v1[0];
|
||||
}
|
||||
|
|
@ -0,0 +1,11 @@
|
|||
void f() {
|
||||
const float4x4 m = float4x4(float4(1.0f, 1.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f));
|
||||
const float4 v1 = m[0];
|
||||
const float a = v1[0];
|
||||
}
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void unused_entry_point() {
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,9 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
void f() {
|
||||
float4x4 const m = float4x4(float4(1.0f, 1.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f));
|
||||
float4 const v1 = m[0];
|
||||
float const a = v1[0];
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,31 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 17
|
||||
; 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 %f "f"
|
||||
%void = OpTypeVoid
|
||||
%1 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v4float = OpTypeVector %float 4
|
||||
%mat4v4float = OpTypeMatrix %v4float 4
|
||||
%float_1 = OpConstant %float 1
|
||||
%11 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
|
||||
%12 = OpConstantComposite %mat4v4float %11 %11 %11 %11
|
||||
%int = OpTypeInt 32 1
|
||||
%int_0 = OpConstant %int 0
|
||||
%unused_entry_point = OpFunction %void None %1
|
||||
%4 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%f = OpFunction %void None %1
|
||||
%6 = OpLabel
|
||||
%15 = OpCompositeExtract %v4float %12 0
|
||||
%16 = OpCompositeExtract %float %15 0
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,5 @@
|
|||
fn f() {
|
||||
let m : mat4x4<f32> = mat4x4<f32>(vec4<f32>(1.0, 1.0, 1.0, 1.0), vec4<f32>(1.0, 1.0, 1.0, 1.0), vec4<f32>(1.0, 1.0, 1.0, 1.0), vec4<f32>(1.0, 1.0, 1.0, 1.0));
|
||||
let v1 : vec4<f32> = m[0];
|
||||
let a : f32 = v1[0];
|
||||
}
|
||||
Binary file not shown.
|
|
@ -0,0 +1,3 @@
|
|||
SKIP: FAILED
|
||||
|
||||
error: unhandled instruction with opcode 150: %712 = OpISubBorrow %710 %107 %470
|
||||
|
|
@ -0,0 +1,3 @@
|
|||
SKIP: FAILED
|
||||
|
||||
error: unhandled instruction with opcode 150: %712 = OpISubBorrow %710 %107 %470
|
||||
|
|
@ -0,0 +1,3 @@
|
|||
SKIP: FAILED
|
||||
|
||||
error: unhandled instruction with opcode 150: %712 = OpISubBorrow %710 %107 %470
|
||||
|
|
@ -0,0 +1,3 @@
|
|||
SKIP: FAILED
|
||||
|
||||
error: unhandled instruction with opcode 150: %712 = OpISubBorrow %710 %107 %470
|
||||
Binary file not shown.
|
|
@ -0,0 +1,10 @@
|
|||
SKIP: FAILED
|
||||
|
||||
../src/ast/array_accessor_expression.cc:29 internal compiler error: TINT_ASSERT(array_)
|
||||
********************************************************************
|
||||
* The tint shader compiler has encountered an unexpected error. *
|
||||
* *
|
||||
* Please help us fix this issue by submitting a bug report at *
|
||||
* crbug.com/tint with the source program that triggered the bug. *
|
||||
********************************************************************
|
||||
|
||||
|
|
@ -0,0 +1,10 @@
|
|||
SKIP: FAILED
|
||||
|
||||
../src/ast/array_accessor_expression.cc:29 internal compiler error: TINT_ASSERT(array_)
|
||||
********************************************************************
|
||||
* The tint shader compiler has encountered an unexpected error. *
|
||||
* *
|
||||
* Please help us fix this issue by submitting a bug report at *
|
||||
* crbug.com/tint with the source program that triggered the bug. *
|
||||
********************************************************************
|
||||
|
||||
|
|
@ -0,0 +1,10 @@
|
|||
SKIP: FAILED
|
||||
|
||||
../src/ast/array_accessor_expression.cc:29 internal compiler error: TINT_ASSERT(array_)
|
||||
********************************************************************
|
||||
* The tint shader compiler has encountered an unexpected error. *
|
||||
* *
|
||||
* Please help us fix this issue by submitting a bug report at *
|
||||
* crbug.com/tint with the source program that triggered the bug. *
|
||||
********************************************************************
|
||||
|
||||
|
|
@ -0,0 +1,10 @@
|
|||
SKIP: FAILED
|
||||
|
||||
../src/ast/array_accessor_expression.cc:29 internal compiler error: TINT_ASSERT(array_)
|
||||
********************************************************************
|
||||
* The tint shader compiler has encountered an unexpected error. *
|
||||
* *
|
||||
* Please help us fix this issue by submitting a bug report at *
|
||||
* crbug.com/tint with the source program that triggered the bug. *
|
||||
********************************************************************
|
||||
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
[[block]] struct Result {
|
||||
values : array<f32>;
|
||||
};
|
||||
|
||||
let width : u32 = 128u;
|
||||
|
||||
[[group(0), binding(0)]] var tex : texture_depth_2d;
|
||||
[[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
|
||||
|
||||
[[stage(compute)]] fn main(
|
||||
[[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>
|
||||
) {
|
||||
result.values[GlobalInvocationId.y * width + GlobalInvocationId.x] = textureLoad(
|
||||
tex, vec2<i32>(i32(GlobalInvocationId.x), i32(GlobalInvocationId.y)), 0);
|
||||
}
|
||||
|
|
@ -0,0 +1,20 @@
|
|||
bug/tint/827.wgsl:8:26 warning: use of deprecated language feature: declare access with var<storage, read_write> instead of using [[access]] decoration
|
||||
[[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
|
||||
^^^
|
||||
|
||||
struct tint_symbol_1 {
|
||||
uint3 GlobalInvocationId : SV_DispatchThreadID;
|
||||
};
|
||||
|
||||
static const uint width = 128u;
|
||||
RWByteAddressBuffer result : register(u1, space0);
|
||||
|
||||
Texture2D tex : register(t0, space0);
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main(tint_symbol_1 tint_symbol) {
|
||||
const uint3 GlobalInvocationId = tint_symbol.GlobalInvocationId;
|
||||
result.Store((4u * ((GlobalInvocationId.y * width) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0), 0)));
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,15 @@
|
|||
SKIP: FAILED
|
||||
|
||||
bug/tint/827.wgsl:8:26 warning: use of deprecated language feature: declare access with var<storage, read_write> instead of using [[access]] decoration
|
||||
[[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
|
||||
^^^
|
||||
|
||||
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:10:76: error: use of undeclared identifier 'tex'
|
||||
result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = tex.read(int2(int(GlobalInvocationId.x), int(GlobalInvocationId.y)), 0);
|
||||
^
|
||||
|
|
@ -0,0 +1,73 @@
|
|||
bug/tint/827.wgsl:8:26 warning: use of deprecated language feature: declare access with var<storage, read_write> instead of using [[access]] decoration
|
||||
[[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
|
||||
^^^
|
||||
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 43
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %width "width"
|
||||
OpName %tex "tex"
|
||||
OpName %Result "Result"
|
||||
OpMemberName %Result 0 "values"
|
||||
OpName %result "result"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %main "main"
|
||||
OpDecorate %tex DescriptorSet 0
|
||||
OpDecorate %tex Binding 0
|
||||
OpDecorate %Result Block
|
||||
OpMemberDecorate %Result 0 Offset 0
|
||||
OpDecorate %_runtimearr_float ArrayStride 4
|
||||
OpDecorate %result DescriptorSet 0
|
||||
OpDecorate %result Binding 1
|
||||
OpDecorate %tint_symbol BuiltIn GlobalInvocationId
|
||||
%uint = OpTypeInt 32 0
|
||||
%width = OpConstant %uint 128
|
||||
%float = OpTypeFloat 32
|
||||
%5 = OpTypeImage %float 2D 1 0 0 1 Unknown
|
||||
%_ptr_UniformConstant_5 = OpTypePointer UniformConstant %5
|
||||
%tex = OpVariable %_ptr_UniformConstant_5 UniformConstant
|
||||
%_runtimearr_float = OpTypeRuntimeArray %float
|
||||
%Result = OpTypeStruct %_runtimearr_float
|
||||
%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
|
||||
%result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||
%void = OpTypeVoid
|
||||
%14 = OpTypeFunction %void
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
|
||||
%v4float = OpTypeVector %float 4
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%main = OpFunction %void None %14
|
||||
%17 = OpLabel
|
||||
%21 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1
|
||||
%22 = OpLoad %uint %21
|
||||
%23 = OpIMul %uint %22 %width
|
||||
%24 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0
|
||||
%25 = OpLoad %uint %24
|
||||
%26 = OpIAdd %uint %23 %25
|
||||
%28 = OpAccessChain %_ptr_StorageBuffer_float %result %uint_0 %26
|
||||
%32 = OpLoad %5 %tex
|
||||
%36 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0
|
||||
%37 = OpLoad %uint %36
|
||||
%35 = OpBitcast %int %37
|
||||
%39 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1
|
||||
%40 = OpLoad %uint %39
|
||||
%38 = OpBitcast %int %40
|
||||
%41 = OpCompositeConstruct %v2int %35 %38
|
||||
%30 = OpImageFetch %v4float %32 %41 Lod %int_0
|
||||
%29 = OpCompositeExtract %float %30 0
|
||||
OpStore %28 %29
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,19 @@
|
|||
bug/tint/827.wgsl:8:26 warning: use of deprecated language feature: declare access with var<storage, read_write> instead of using [[access]] decoration
|
||||
[[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
|
||||
^^^
|
||||
|
||||
[[block]]
|
||||
struct Result {
|
||||
values : array<f32>;
|
||||
};
|
||||
|
||||
let width : u32 = 128u;
|
||||
|
||||
[[group(0), binding(0)]] var tex : texture_depth_2d;
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read_write> result : Result;
|
||||
|
||||
[[stage(compute)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>) {
|
||||
result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = textureLoad(tex, vec2<i32>(i32(GlobalInvocationId.x), i32(GlobalInvocationId.y)), 0);
|
||||
}
|
||||
Loading…
Reference in New Issue