test: Add missing .expected files
Theses were missing from https://dawn-review.googlesource.com/c/tint/+/56140 Change-Id: Iece1c69923024da2ce9473e205c18eb9568872b9 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56548 Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Reviewed-by: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
e58ecaa7cb
commit
5e2d8af6ad
|
|
@ -0,0 +1,67 @@
|
|||
void Set_uint4(inout uint4 vec, int idx, uint val) {
|
||||
switch(idx) {
|
||||
case 0: vec[0] = val; break;
|
||||
case 1: vec[1] = val; break;
|
||||
case 2: vec[2] = val; break;
|
||||
case 3: vec[3] = val; break;
|
||||
}
|
||||
}
|
||||
|
||||
Texture2D<float4> src : register(t0, space0);
|
||||
Texture2D<float4> tint_symbol : register(t1, space0);
|
||||
RWByteAddressBuffer output : register(u2, space0);
|
||||
cbuffer cbuffer_uniforms : register(b3, space0) {
|
||||
uint4 uniforms[1];
|
||||
};
|
||||
|
||||
uint ConvertToFp16FloatValue(float fp32) {
|
||||
return 1u;
|
||||
}
|
||||
|
||||
struct tint_symbol_2 {
|
||||
uint3 GlobalInvocationID : SV_DispatchThreadID;
|
||||
};
|
||||
|
||||
[numthreads(1, 1, 1)]
|
||||
void main(tint_symbol_2 tint_symbol_1) {
|
||||
const uint3 GlobalInvocationID = tint_symbol_1.GlobalInvocationID;
|
||||
int2 tint_tmp;
|
||||
src.GetDimensions(tint_tmp.x, tint_tmp.y);
|
||||
int2 size = tint_tmp;
|
||||
int2 dstTexCoord = int2(GlobalInvocationID.xy);
|
||||
int2 srcTexCoord = dstTexCoord;
|
||||
const int scalar_offset = (0u) / 4;
|
||||
if ((uniforms[scalar_offset / 4][scalar_offset % 4] == 1u)) {
|
||||
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
|
||||
}
|
||||
float4 srcColor = src.Load(int3(srcTexCoord, 0));
|
||||
float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0));
|
||||
bool success = true;
|
||||
uint4 srcColorBits = uint4(0u, 0u, 0u, 0u);
|
||||
uint4 dstColorBits = uint4(dstColor);
|
||||
{
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
const int scalar_offset_1 = (12u) / 4;
|
||||
if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) {
|
||||
break;
|
||||
}
|
||||
Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
|
||||
bool tint_tmp_1 = success;
|
||||
if (tint_tmp_1) {
|
||||
tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
|
||||
}
|
||||
success = (tint_tmp_1);
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
output.Store((4u * outputIndex), asuint(uint(1)));
|
||||
} else {
|
||||
output.Store((4u * outputIndex), asuint(uint(0)));
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
@ -0,0 +1,51 @@
|
|||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint dstTextureFlipY;
|
||||
/* 0x0004 */ uint isFloat16;
|
||||
/* 0x0008 */ uint isRGB10A2Unorm;
|
||||
/* 0x000c */ uint channelCount;
|
||||
};
|
||||
struct OutputBuf {
|
||||
/* 0x0000 */ uint result[1];
|
||||
};
|
||||
|
||||
uint ConvertToFp16FloatValue(float fp32) {
|
||||
return 1u;
|
||||
}
|
||||
|
||||
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture(0)]], texture2d<float, access::sample> tint_symbol_3 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
|
||||
int2 size = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
|
||||
int2 dstTexCoord = int2(GlobalInvocationID.xy);
|
||||
int2 srcTexCoord = dstTexCoord;
|
||||
if ((uniforms.dstTextureFlipY == 1u)) {
|
||||
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
|
||||
}
|
||||
float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
|
||||
float4 dstColor = tint_symbol_3.read(uint2(dstTexCoord), 0);
|
||||
bool success = true;
|
||||
uint4 srcColorBits = 0u;
|
||||
uint4 dstColorBits = uint4(dstColor);
|
||||
{
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
if (!((i < uniforms.channelCount))) {
|
||||
break;
|
||||
}
|
||||
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
|
||||
success = (success && (srcColorBits[i] == dstColorBits[i]));
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
output.result[outputIndex] = uint(1);
|
||||
} else {
|
||||
output.result[outputIndex] = uint(0);
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,224 @@
|
|||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 135
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability ImageQuery
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %tint_symbol
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpName %src "src"
|
||||
OpName %dst "dst"
|
||||
OpName %OutputBuf "OutputBuf"
|
||||
OpMemberName %OutputBuf 0 "result"
|
||||
OpName %output "output"
|
||||
OpName %Uniforms "Uniforms"
|
||||
OpMemberName %Uniforms 0 "dstTextureFlipY"
|
||||
OpMemberName %Uniforms 1 "isFloat16"
|
||||
OpMemberName %Uniforms 2 "isRGB10A2Unorm"
|
||||
OpMemberName %Uniforms 3 "channelCount"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %ConvertToFp16FloatValue "ConvertToFp16FloatValue"
|
||||
OpName %fp32 "fp32"
|
||||
OpName %main "main"
|
||||
OpName %size "size"
|
||||
OpName %dstTexCoord "dstTexCoord"
|
||||
OpName %srcTexCoord "srcTexCoord"
|
||||
OpName %srcColor "srcColor"
|
||||
OpName %dstColor "dstColor"
|
||||
OpName %success "success"
|
||||
OpName %srcColorBits "srcColorBits"
|
||||
OpName %dstColorBits "dstColorBits"
|
||||
OpName %i "i"
|
||||
OpName %outputIndex "outputIndex"
|
||||
OpDecorate %src DescriptorSet 0
|
||||
OpDecorate %src Binding 0
|
||||
OpDecorate %dst DescriptorSet 0
|
||||
OpDecorate %dst Binding 1
|
||||
OpDecorate %OutputBuf Block
|
||||
OpMemberDecorate %OutputBuf 0 Offset 0
|
||||
OpDecorate %_runtimearr_uint ArrayStride 4
|
||||
OpDecorate %output DescriptorSet 0
|
||||
OpDecorate %output Binding 2
|
||||
OpDecorate %Uniforms Block
|
||||
OpMemberDecorate %Uniforms 0 Offset 0
|
||||
OpMemberDecorate %Uniforms 1 Offset 4
|
||||
OpMemberDecorate %Uniforms 2 Offset 8
|
||||
OpMemberDecorate %Uniforms 3 Offset 12
|
||||
OpDecorate %uniforms NonWritable
|
||||
OpDecorate %uniforms DescriptorSet 0
|
||||
OpDecorate %uniforms Binding 3
|
||||
OpDecorate %tint_symbol BuiltIn GlobalInvocationId
|
||||
%float = OpTypeFloat 32
|
||||
%3 = OpTypeImage %float 2D 0 0 0 1 Unknown
|
||||
%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
|
||||
%src = OpVariable %_ptr_UniformConstant_3 UniformConstant
|
||||
%dst = OpVariable %_ptr_UniformConstant_3 UniformConstant
|
||||
%uint = OpTypeInt 32 0
|
||||
%_runtimearr_uint = OpTypeRuntimeArray %uint
|
||||
%OutputBuf = OpTypeStruct %_runtimearr_uint
|
||||
%_ptr_StorageBuffer_OutputBuf = OpTypePointer StorageBuffer %OutputBuf
|
||||
%output = OpVariable %_ptr_StorageBuffer_OutputBuf StorageBuffer
|
||||
%Uniforms = OpTypeStruct %uint %uint %uint %uint
|
||||
%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
|
||||
%uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%tint_symbol = OpVariable %_ptr_Input_v3uint Input
|
||||
%17 = OpTypeFunction %uint %float
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%void = OpTypeVoid
|
||||
%22 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%_ptr_Function_v2int = OpTypePointer Function %v2int
|
||||
%33 = OpConstantNull %v2int
|
||||
%v2uint = OpTypeVector %uint 2
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%int_1 = OpConstant %int 1
|
||||
%v4float = OpTypeVector %float 4
|
||||
%_ptr_Function_v4float = OpTypePointer Function %v4float
|
||||
%64 = OpConstantNull %v4float
|
||||
%true = OpConstantTrue %bool
|
||||
%_ptr_Function_bool = OpTypePointer Function %bool
|
||||
%72 = OpConstantNull %bool
|
||||
%v4uint = OpTypeVector %uint 4
|
||||
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
|
||||
%76 = OpConstantNull %v4uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%82 = OpConstantNull %uint
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_ptr_Function_float = OpTypePointer Function %float
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
|
||||
%ConvertToFp16FloatValue = OpFunction %uint None %17
|
||||
%fp32 = OpFunctionParameter %float
|
||||
%20 = OpLabel
|
||||
OpReturnValue %uint_1
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %22
|
||||
%25 = OpLabel
|
||||
%size = OpVariable %_ptr_Function_v2int Function %33
|
||||
%dstTexCoord = OpVariable %_ptr_Function_v2int Function %33
|
||||
%srcTexCoord = OpVariable %_ptr_Function_v2int Function %33
|
||||
%srcColor = OpVariable %_ptr_Function_v4float Function %64
|
||||
%dstColor = OpVariable %_ptr_Function_v4float Function %64
|
||||
%success = OpVariable %_ptr_Function_bool Function %72
|
||||
%srcColorBits = OpVariable %_ptr_Function_v4uint Function %76
|
||||
%dstColorBits = OpVariable %_ptr_Function_v4uint Function %76
|
||||
%i = OpVariable %_ptr_Function_uint Function %82
|
||||
%outputIndex = OpVariable %_ptr_Function_uint Function %82
|
||||
%29 = OpLoad %3 %src
|
||||
%26 = OpImageQuerySizeLod %v2int %29 %int_0
|
||||
OpStore %size %26
|
||||
%36 = OpLoad %v3uint %tint_symbol
|
||||
%37 = OpVectorShuffle %v2uint %36 %36 0 1
|
||||
%34 = OpBitcast %v2int %37
|
||||
OpStore %dstTexCoord %34
|
||||
%39 = OpLoad %v2int %dstTexCoord
|
||||
OpStore %srcTexCoord %39
|
||||
%43 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%44 = OpLoad %uint %43
|
||||
%45 = OpIEqual %bool %44 %uint_1
|
||||
OpSelectionMerge %47 None
|
||||
OpBranchConditional %45 %48 %47
|
||||
%48 = OpLabel
|
||||
%50 = OpAccessChain %_ptr_Function_int %srcTexCoord %uint_1
|
||||
%51 = OpAccessChain %_ptr_Function_int %size %uint_1
|
||||
%52 = OpLoad %int %51
|
||||
%53 = OpAccessChain %_ptr_Function_int %dstTexCoord %uint_1
|
||||
%54 = OpLoad %int %53
|
||||
%55 = OpISub %int %52 %54
|
||||
%57 = OpISub %int %55 %int_1
|
||||
OpStore %50 %57
|
||||
OpBranch %47
|
||||
%47 = OpLabel
|
||||
%60 = OpLoad %3 %src
|
||||
%61 = OpLoad %v2int %srcTexCoord
|
||||
%58 = OpImageFetch %v4float %60 %61 Lod %int_0
|
||||
OpStore %srcColor %58
|
||||
%66 = OpLoad %3 %dst
|
||||
%67 = OpLoad %v2int %dstTexCoord
|
||||
%65 = OpImageFetch %v4float %66 %67 Lod %int_0
|
||||
OpStore %dstColor %65
|
||||
OpStore %success %true
|
||||
%78 = OpLoad %v4float %dstColor
|
||||
%77 = OpConvertFToU %v4uint %78
|
||||
OpStore %dstColorBits %77
|
||||
OpStore %i %uint_0
|
||||
OpBranch %83
|
||||
%83 = OpLabel
|
||||
OpLoopMerge %84 %85 None
|
||||
OpBranch %86
|
||||
%86 = OpLabel
|
||||
%88 = OpLoad %uint %i
|
||||
%90 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3
|
||||
%91 = OpLoad %uint %90
|
||||
%92 = OpULessThan %bool %88 %91
|
||||
%87 = OpLogicalNot %bool %92
|
||||
OpSelectionMerge %93 None
|
||||
OpBranchConditional %87 %94 %93
|
||||
%94 = OpLabel
|
||||
OpBranch %84
|
||||
%93 = OpLabel
|
||||
%95 = OpLoad %uint %i
|
||||
%96 = OpAccessChain %_ptr_Function_uint %srcColorBits %95
|
||||
%98 = OpLoad %uint %i
|
||||
%100 = OpAccessChain %_ptr_Function_float %srcColor %98
|
||||
%101 = OpLoad %float %100
|
||||
%97 = OpFunctionCall %uint %ConvertToFp16FloatValue %101
|
||||
OpStore %96 %97
|
||||
%102 = OpLoad %bool %success
|
||||
OpSelectionMerge %103 None
|
||||
OpBranchConditional %102 %104 %103
|
||||
%104 = OpLabel
|
||||
%105 = OpLoad %uint %i
|
||||
%106 = OpAccessChain %_ptr_Function_uint %srcColorBits %105
|
||||
%107 = OpLoad %uint %106
|
||||
%108 = OpLoad %uint %i
|
||||
%109 = OpAccessChain %_ptr_Function_uint %dstColorBits %108
|
||||
%110 = OpLoad %uint %109
|
||||
%111 = OpIEqual %bool %107 %110
|
||||
OpBranch %103
|
||||
%103 = OpLabel
|
||||
%112 = OpPhi %bool %102 %93 %111 %104
|
||||
OpStore %success %112
|
||||
OpBranch %85
|
||||
%85 = OpLabel
|
||||
%113 = OpLoad %uint %i
|
||||
%114 = OpIAdd %uint %113 %uint_1
|
||||
OpStore %i %114
|
||||
OpBranch %83
|
||||
%84 = OpLabel
|
||||
%116 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1
|
||||
%117 = OpLoad %uint %116
|
||||
%119 = OpAccessChain %_ptr_Function_int %size %uint_0
|
||||
%120 = OpLoad %int %119
|
||||
%118 = OpBitcast %uint %120
|
||||
%121 = OpIMul %uint %117 %118
|
||||
%122 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0
|
||||
%123 = OpLoad %uint %122
|
||||
%124 = OpIAdd %uint %121 %123
|
||||
OpStore %outputIndex %124
|
||||
%126 = OpLoad %bool %success
|
||||
OpSelectionMerge %127 None
|
||||
OpBranchConditional %126 %128 %129
|
||||
%128 = OpLabel
|
||||
%130 = OpLoad %uint %outputIndex
|
||||
%132 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %130
|
||||
OpStore %132 %uint_1
|
||||
OpBranch %127
|
||||
%129 = OpLabel
|
||||
%133 = OpLoad %uint %outputIndex
|
||||
%134 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %133
|
||||
OpStore %134 %uint_0
|
||||
OpBranch %127
|
||||
%127 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
|
|
@ -0,0 +1,59 @@
|
|||
[[block]]
|
||||
struct Uniforms {
|
||||
dstTextureFlipY : u32;
|
||||
isFloat16 : u32;
|
||||
isRGB10A2Unorm : u32;
|
||||
channelCount : u32;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct OutputBuf {
|
||||
result : [[stride(4)]] array<u32>;
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var src : texture_2d<f32>;
|
||||
|
||||
[[group(0), binding(1)]] var dst : texture_2d<f32>;
|
||||
|
||||
[[group(0), binding(2)]] var<storage, read_write> output : OutputBuf;
|
||||
|
||||
[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
|
||||
|
||||
fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
|
||||
return 1u;
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(1, 1, 1)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
var size : vec2<i32> = textureDimensions(src);
|
||||
var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
|
||||
var srcTexCoord : vec2<i32> = dstTexCoord;
|
||||
if ((uniforms.dstTextureFlipY == 1u)) {
|
||||
srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
|
||||
}
|
||||
var srcColor : vec4<f32> = textureLoad(src, srcTexCoord, 0);
|
||||
var dstColor : vec4<f32> = textureLoad(dst, dstTexCoord, 0);
|
||||
var success : bool = true;
|
||||
var srcColorBits : vec4<u32>;
|
||||
var dstColorBits : vec4<u32> = vec4<u32>(dstColor);
|
||||
{
|
||||
var i : u32 = 0u;
|
||||
loop {
|
||||
if (!((i < uniforms.channelCount))) {
|
||||
break;
|
||||
}
|
||||
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
|
||||
success = (success && (srcColorBits[i] == dstColorBits[i]));
|
||||
|
||||
continuing {
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
output.result[outputIndex] = u32(1);
|
||||
} else {
|
||||
output.result[outputIndex] = u32(0);
|
||||
}
|
||||
}
|
||||
Loading…
Reference in New Issue