mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-07-14 17:16:01 +00:00
writer/hlsl: Add missing parenthesis around UBO ternary op
Bug: tint:913 Change-Id: I2edbab363cb03e6ce64b5c6bddf184bf92438521 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55340 Auto-Submit: Ben Clayton <bclayton@google.com> Reviewed-by: Sarah Mashayekhi <sarahmashay@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
5eadb83a3f
commit
0f916164ae
@ -564,8 +564,8 @@ bool GeneratorImpl::EmitUniformBufferAccess(
|
||||
make_indent(ss << ";" << std::endl);
|
||||
pre << ss.str();
|
||||
|
||||
out << "(" << scalar_offset << " & 2) ? " << ubo_load
|
||||
<< ".zw : " << ubo_load << ".xy";
|
||||
out << "((" << scalar_offset << " & 2) ? " << ubo_load
|
||||
<< ".zw : " << ubo_load << ".xy)";
|
||||
return true;
|
||||
};
|
||||
// vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
|
||||
|
||||
@ -868,7 +868,7 @@ void main_1() {
|
||||
uv = x_762;
|
||||
const int scalar_offset = (0u) / 4;
|
||||
uint4 ubo_load = x_188[scalar_offset / 4];
|
||||
const float2 x_191 = asfloat((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy);
|
||||
const float2 x_191 = asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy));
|
||||
const QuicksortObject x_763 = obj;
|
||||
const tint_array_wrapper tint_symbol_51 = {{0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
|
||||
const QuicksortObject tint_symbol_52 = {tint_symbol_51};
|
||||
|
||||
70
test/bug/tint/913.wgsl
Normal file
70
test/bug/tint/913.wgsl
Normal file
@ -0,0 +1,70 @@
|
||||
[[block]] struct Uniforms {
|
||||
dstTextureFlipY : u32;
|
||||
channelCount : u32;
|
||||
srcCopyOrigin : vec2<u32>;
|
||||
dstCopyOrigin : vec2<u32>;
|
||||
copySize : vec2<u32>;
|
||||
};
|
||||
[[block]] struct OutputBuf {
|
||||
result : 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 aboutEqual(value : f32, expect : f32) -> bool {
|
||||
// The value diff should be smaller than the hard coded tolerance.
|
||||
return abs(value - expect) < 0.001;
|
||||
}
|
||||
[[stage(compute), workgroup_size(1, 1, 1)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
let srcSize : vec2<i32> = textureDimensions(src);
|
||||
let dstSize : vec2<i32> = textureDimensions(dst);
|
||||
let dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
|
||||
let nonCoveredColor : vec4<f32> =
|
||||
vec4<f32>(0.0, 1.0, 0.0, 1.0); // should be green
|
||||
|
||||
var success : bool = true;
|
||||
if (dstTexCoord.x < uniforms.dstCopyOrigin.x ||
|
||||
dstTexCoord.y < uniforms.dstCopyOrigin.y ||
|
||||
dstTexCoord.x >= uniforms.dstCopyOrigin.x + uniforms.copySize.x ||
|
||||
dstTexCoord.y >= uniforms.dstCopyOrigin.y + uniforms.copySize.y) {
|
||||
success = success &&
|
||||
all(textureLoad(dst, vec2<i32>(dstTexCoord), 0) == nonCoveredColor);
|
||||
} else {
|
||||
// Calculate source texture coord.
|
||||
var srcTexCoord : vec2<u32> = dstTexCoord - uniforms.dstCopyOrigin +
|
||||
uniforms.srcCopyOrigin;
|
||||
// Note that |flipY| equals flip src texture firstly and then do copy from src
|
||||
// subrect to dst subrect. This helps on blink part to handle some input texture
|
||||
// which is flipped and need to unpack flip during the copy.
|
||||
// We need to calculate the expect y coord based on this rule.
|
||||
if (uniforms.dstTextureFlipY == 1u) {
|
||||
srcTexCoord.y = u32(srcSize.y) - srcTexCoord.y - 1u;
|
||||
}
|
||||
|
||||
let srcColor : vec4<f32> = textureLoad(src, vec2<i32>(srcTexCoord), 0);
|
||||
let dstColor : vec4<f32> = textureLoad(dst, vec2<i32>(dstTexCoord), 0);
|
||||
|
||||
// Not use loop and variable index format to workaround
|
||||
// crbug.com/tint/638.
|
||||
if (uniforms.channelCount == 2u) { // All have rg components.
|
||||
success = success &&
|
||||
aboutEqual(dstColor.r, srcColor.r) &&
|
||||
aboutEqual(dstColor.g, srcColor.g);
|
||||
} else {
|
||||
success = success &&
|
||||
aboutEqual(dstColor.r, srcColor.r) &&
|
||||
aboutEqual(dstColor.g, srcColor.g) &&
|
||||
aboutEqual(dstColor.b, srcColor.b) &&
|
||||
aboutEqual(dstColor.a, srcColor.a);
|
||||
}
|
||||
}
|
||||
let outputIndex : u32 = GlobalInvocationID.y * u32(dstSize.x) +
|
||||
GlobalInvocationID.x;
|
||||
if (success) {
|
||||
output.result[outputIndex] = 1u;
|
||||
} else {
|
||||
output.result[outputIndex] = 0u;
|
||||
}
|
||||
}
|
||||
102
test/bug/tint/913.wgsl.expected.hlsl
Normal file
102
test/bug/tint/913.wgsl.expected.hlsl
Normal file
@ -0,0 +1,102 @@
|
||||
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[2];
|
||||
};
|
||||
|
||||
bool aboutEqual(float value, float expect) {
|
||||
return (abs((value - expect)) < 0.001f);
|
||||
}
|
||||
|
||||
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);
|
||||
const int2 srcSize = tint_tmp;
|
||||
int2 tint_tmp_1;
|
||||
tint_symbol.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
|
||||
const int2 dstSize = tint_tmp_1;
|
||||
const uint2 dstTexCoord = uint2(GlobalInvocationID.xy);
|
||||
const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
|
||||
bool success = true;
|
||||
const int scalar_offset = (16u) / 4;
|
||||
bool tint_tmp_2 = (dstTexCoord.x < uniforms[scalar_offset / 4][scalar_offset % 4]);
|
||||
if (!tint_tmp_2) {
|
||||
const int scalar_offset_1 = (20u) / 4;
|
||||
tint_tmp_2 = (dstTexCoord.y < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
|
||||
}
|
||||
bool tint_tmp_3 = (tint_tmp_2);
|
||||
if (!tint_tmp_3) {
|
||||
const int scalar_offset_2 = (16u) / 4;
|
||||
const int scalar_offset_3 = (24u) / 4;
|
||||
tint_tmp_3 = (dstTexCoord.x >= (uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4] + uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]));
|
||||
}
|
||||
bool tint_tmp_4 = (tint_tmp_3);
|
||||
if (!tint_tmp_4) {
|
||||
const int scalar_offset_4 = (20u) / 4;
|
||||
const int scalar_offset_5 = (28u) / 4;
|
||||
tint_tmp_4 = (dstTexCoord.y >= (uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4] + uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]));
|
||||
}
|
||||
if ((tint_tmp_4)) {
|
||||
bool tint_tmp_5 = success;
|
||||
if (tint_tmp_5) {
|
||||
tint_tmp_5 = all((tint_symbol.Load(int3(dstTexCoord, 0), 0) == nonCoveredColor));
|
||||
}
|
||||
success = (tint_tmp_5);
|
||||
} else {
|
||||
const int scalar_offset_6 = (16u) / 4;
|
||||
uint4 ubo_load = uniforms[scalar_offset_6 / 4];
|
||||
const int scalar_offset_7 = (8u) / 4;
|
||||
uint4 ubo_load_1 = uniforms[scalar_offset_7 / 4];
|
||||
uint2 srcTexCoord = ((dstTexCoord - ((scalar_offset_6 & 2) ? ubo_load.zw : ubo_load.xy)) + ((scalar_offset_7 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
|
||||
const int scalar_offset_8 = (0u) / 4;
|
||||
if ((uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4] == 1u)) {
|
||||
srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
|
||||
}
|
||||
const float4 srcColor = src.Load(int3(srcTexCoord, 0), 0);
|
||||
const float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0), 0);
|
||||
const int scalar_offset_9 = (4u) / 4;
|
||||
if ((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] == 2u)) {
|
||||
bool tint_tmp_6 = success;
|
||||
if (tint_tmp_6) {
|
||||
tint_tmp_6 = aboutEqual(dstColor.r, srcColor.r);
|
||||
}
|
||||
bool tint_tmp_7 = (tint_tmp_6);
|
||||
if (tint_tmp_7) {
|
||||
tint_tmp_7 = aboutEqual(dstColor.g, srcColor.g);
|
||||
}
|
||||
success = (tint_tmp_7);
|
||||
} else {
|
||||
bool tint_tmp_8 = success;
|
||||
if (tint_tmp_8) {
|
||||
tint_tmp_8 = aboutEqual(dstColor.r, srcColor.r);
|
||||
}
|
||||
bool tint_tmp_9 = (tint_tmp_8);
|
||||
if (tint_tmp_9) {
|
||||
tint_tmp_9 = aboutEqual(dstColor.g, srcColor.g);
|
||||
}
|
||||
bool tint_tmp_10 = (tint_tmp_9);
|
||||
if (tint_tmp_10) {
|
||||
tint_tmp_10 = aboutEqual(dstColor.b, srcColor.b);
|
||||
}
|
||||
bool tint_tmp_11 = (tint_tmp_10);
|
||||
if (tint_tmp_11) {
|
||||
tint_tmp_11 = aboutEqual(dstColor.a, srcColor.a);
|
||||
}
|
||||
success = (tint_tmp_11);
|
||||
}
|
||||
}
|
||||
const uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
output.Store((4u * outputIndex), asuint(1u));
|
||||
} else {
|
||||
output.Store((4u * outputIndex), asuint(0u));
|
||||
}
|
||||
return;
|
||||
}
|
||||
48
test/bug/tint/913.wgsl.expected.msl
Normal file
48
test/bug/tint/913.wgsl.expected.msl
Normal file
@ -0,0 +1,48 @@
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint dstTextureFlipY;
|
||||
/* 0x0004 */ uint channelCount;
|
||||
/* 0x0008 */ packed_uint2 srcCopyOrigin;
|
||||
/* 0x0010 */ packed_uint2 dstCopyOrigin;
|
||||
/* 0x0018 */ packed_uint2 copySize;
|
||||
};
|
||||
struct OutputBuf {
|
||||
/* 0x0000 */ uint result[1];
|
||||
};
|
||||
|
||||
bool aboutEqual(float value, float expect) {
|
||||
return ( fabs((value - expect)) < 0.001f);
|
||||
}
|
||||
|
||||
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 const srcSize = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
|
||||
int2 const dstSize = int2(tint_symbol_3.get_width(), tint_symbol_3.get_height());
|
||||
uint2 const dstTexCoord = uint2(GlobalInvocationID.xy);
|
||||
float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
|
||||
bool success = true;
|
||||
if (((((dstTexCoord.x < uniforms.dstCopyOrigin.x) || (dstTexCoord.y < uniforms.dstCopyOrigin.y)) || (dstTexCoord.x >= (uniforms.dstCopyOrigin.x + uniforms.copySize.x))) || (dstTexCoord.y >= (uniforms.dstCopyOrigin.y + uniforms.copySize.y)))) {
|
||||
success = (success && all((tint_symbol_3.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
|
||||
} else {
|
||||
uint2 srcTexCoord = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin);
|
||||
if ((uniforms.dstTextureFlipY == 1u)) {
|
||||
srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
|
||||
}
|
||||
float4 const srcColor = tint_symbol_2.read(uint2(int2(srcTexCoord)), 0);
|
||||
float4 const dstColor = tint_symbol_3.read(uint2(int2(dstTexCoord)), 0);
|
||||
if ((uniforms.channelCount == 2u)) {
|
||||
success = ((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g));
|
||||
} else {
|
||||
success = ((((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)) && aboutEqual(dstColor.b, srcColor.b)) && aboutEqual(dstColor.a, srcColor.a));
|
||||
}
|
||||
}
|
||||
uint const outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
output.result[outputIndex] = 1u;
|
||||
} else {
|
||||
output.result[outputIndex] = 0u;
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
294
test/bug/tint/913.wgsl.expected.spvasm
Normal file
294
test/bug/tint/913.wgsl.expected.spvasm
Normal file
@ -0,0 +1,294 @@
|
||||
; SPIR-V
|
||||
; Version: 1.3
|
||||
; Generator: Google Tint Compiler; 0
|
||||
; Bound: 191
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability ImageQuery
|
||||
%25 = OpExtInstImport "GLSL.std.450"
|
||||
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 "channelCount"
|
||||
OpMemberName %Uniforms 2 "srcCopyOrigin"
|
||||
OpMemberName %Uniforms 3 "dstCopyOrigin"
|
||||
OpMemberName %Uniforms 4 "copySize"
|
||||
OpName %uniforms "uniforms"
|
||||
OpName %tint_symbol "tint_symbol"
|
||||
OpName %aboutEqual "aboutEqual"
|
||||
OpName %value "value"
|
||||
OpName %expect "expect"
|
||||
OpName %main "main"
|
||||
OpName %success "success"
|
||||
OpName %srcTexCoord "srcTexCoord"
|
||||
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 16
|
||||
OpMemberDecorate %Uniforms 4 Offset 24
|
||||
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
|
||||
%v2uint = OpTypeVector %uint 2
|
||||
%Uniforms = OpTypeStruct %uint %uint %v2uint %v2uint %v2uint
|
||||
%_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
|
||||
%bool = OpTypeBool
|
||||
%18 = OpTypeFunction %bool %float %float
|
||||
%float_0_00100000005 = OpConstant %float 0.00100000005
|
||||
%void = OpTypeVoid
|
||||
%29 = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%v2int = OpTypeVector %int 2
|
||||
%int_0 = OpConstant %int 0
|
||||
%v4float = OpTypeVector %float 4
|
||||
%float_0 = OpConstant %float 0
|
||||
%float_1 = OpConstant %float 1
|
||||
%46 = OpConstantComposite %v4float %float_0 %float_1 %float_0 %float_1
|
||||
%true = OpConstantTrue %bool
|
||||
%_ptr_Function_bool = OpTypePointer Function %bool
|
||||
%50 = OpConstantNull %bool
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%v4bool = OpTypeVector %bool 4
|
||||
%_ptr_Uniform_v2uint = OpTypePointer Uniform %v2uint
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%_ptr_Function_v2uint = OpTypePointer Function %v2uint
|
||||
%110 = OpConstantNull %v2uint
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
|
||||
%aboutEqual = OpFunction %bool None %18
|
||||
%value = OpFunctionParameter %float
|
||||
%expect = OpFunctionParameter %float
|
||||
%23 = OpLabel
|
||||
%26 = OpFSub %float %value %expect
|
||||
%24 = OpExtInst %float %25 FAbs %26
|
||||
%28 = OpFOrdLessThan %bool %24 %float_0_00100000005
|
||||
OpReturnValue %28
|
||||
OpFunctionEnd
|
||||
%main = OpFunction %void None %29
|
||||
%32 = OpLabel
|
||||
%success = OpVariable %_ptr_Function_bool Function %50
|
||||
%srcTexCoord = OpVariable %_ptr_Function_v2uint Function %110
|
||||
%36 = OpLoad %3 %src
|
||||
%33 = OpImageQuerySizeLod %v2int %36 %int_0
|
||||
%39 = OpLoad %3 %dst
|
||||
%38 = OpImageQuerySizeLod %v2int %39 %int_0
|
||||
%41 = OpLoad %v3uint %tint_symbol
|
||||
%42 = OpVectorShuffle %v2uint %41 %41 0 1
|
||||
OpStore %success %true
|
||||
%51 = OpCompositeExtract %uint %42 0
|
||||
%55 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0
|
||||
%56 = OpLoad %uint %55
|
||||
%57 = OpULessThan %bool %51 %56
|
||||
OpSelectionMerge %58 None
|
||||
OpBranchConditional %57 %58 %59
|
||||
%59 = OpLabel
|
||||
%60 = OpCompositeExtract %uint %42 1
|
||||
%62 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1
|
||||
%63 = OpLoad %uint %62
|
||||
%64 = OpULessThan %bool %60 %63
|
||||
OpBranch %58
|
||||
%58 = OpLabel
|
||||
%65 = OpPhi %bool %57 %32 %64 %59
|
||||
OpSelectionMerge %66 None
|
||||
OpBranchConditional %65 %66 %67
|
||||
%67 = OpLabel
|
||||
%68 = OpCompositeExtract %uint %42 0
|
||||
%69 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_0
|
||||
%70 = OpLoad %uint %69
|
||||
%72 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_0
|
||||
%73 = OpLoad %uint %72
|
||||
%74 = OpIAdd %uint %70 %73
|
||||
%75 = OpUGreaterThanEqual %bool %68 %74
|
||||
OpBranch %66
|
||||
%66 = OpLabel
|
||||
%76 = OpPhi %bool %65 %58 %75 %67
|
||||
OpSelectionMerge %77 None
|
||||
OpBranchConditional %76 %77 %78
|
||||
%78 = OpLabel
|
||||
%79 = OpCompositeExtract %uint %42 1
|
||||
%80 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_3 %uint_1
|
||||
%81 = OpLoad %uint %80
|
||||
%82 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_4 %uint_1
|
||||
%83 = OpLoad %uint %82
|
||||
%84 = OpIAdd %uint %81 %83
|
||||
%85 = OpUGreaterThanEqual %bool %79 %84
|
||||
OpBranch %77
|
||||
%77 = OpLabel
|
||||
%86 = OpPhi %bool %76 %66 %85 %78
|
||||
OpSelectionMerge %87 None
|
||||
OpBranchConditional %86 %88 %89
|
||||
%88 = OpLabel
|
||||
%90 = OpLoad %bool %success
|
||||
OpSelectionMerge %91 None
|
||||
OpBranchConditional %90 %92 %91
|
||||
%92 = OpLabel
|
||||
%95 = OpLoad %3 %dst
|
||||
%96 = OpBitcast %v2int %42
|
||||
%94 = OpImageFetch %v4float %95 %96 Lod %int_0
|
||||
%97 = OpFOrdEqual %v4bool %94 %46
|
||||
%93 = OpAll %bool %97
|
||||
OpBranch %91
|
||||
%91 = OpLabel
|
||||
%99 = OpPhi %bool %90 %88 %93 %92
|
||||
OpStore %success %99
|
||||
OpBranch %87
|
||||
%89 = OpLabel
|
||||
%101 = OpAccessChain %_ptr_Uniform_v2uint %uniforms %uint_3
|
||||
%102 = OpLoad %v2uint %101
|
||||
%103 = OpISub %v2uint %42 %102
|
||||
%105 = OpAccessChain %_ptr_Uniform_v2uint %uniforms %uint_2
|
||||
%106 = OpLoad %v2uint %105
|
||||
%107 = OpIAdd %v2uint %103 %106
|
||||
OpStore %srcTexCoord %107
|
||||
%111 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
|
||||
%112 = OpLoad %uint %111
|
||||
%113 = OpIEqual %bool %112 %uint_1
|
||||
OpSelectionMerge %114 None
|
||||
OpBranchConditional %113 %115 %114
|
||||
%115 = OpLabel
|
||||
%117 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1
|
||||
%119 = OpCompositeExtract %int %33 1
|
||||
%118 = OpBitcast %uint %119
|
||||
%120 = OpAccessChain %_ptr_Function_uint %srcTexCoord %uint_1
|
||||
%121 = OpLoad %uint %120
|
||||
%122 = OpISub %uint %118 %121
|
||||
%123 = OpISub %uint %122 %uint_1
|
||||
OpStore %117 %123
|
||||
OpBranch %114
|
||||
%114 = OpLabel
|
||||
%125 = OpLoad %3 %src
|
||||
%127 = OpLoad %v2uint %srcTexCoord
|
||||
%126 = OpBitcast %v2int %127
|
||||
%124 = OpImageFetch %v4float %125 %126 Lod %int_0
|
||||
%129 = OpLoad %3 %dst
|
||||
%130 = OpBitcast %v2int %42
|
||||
%128 = OpImageFetch %v4float %129 %130 Lod %int_0
|
||||
%131 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
|
||||
%132 = OpLoad %uint %131
|
||||
%133 = OpIEqual %bool %132 %uint_2
|
||||
OpSelectionMerge %134 None
|
||||
OpBranchConditional %133 %135 %136
|
||||
%135 = OpLabel
|
||||
%137 = OpLoad %bool %success
|
||||
OpSelectionMerge %138 None
|
||||
OpBranchConditional %137 %139 %138
|
||||
%139 = OpLabel
|
||||
%141 = OpCompositeExtract %float %128 0
|
||||
%142 = OpCompositeExtract %float %124 0
|
||||
%140 = OpFunctionCall %bool %aboutEqual %141 %142
|
||||
OpBranch %138
|
||||
%138 = OpLabel
|
||||
%143 = OpPhi %bool %137 %135 %140 %139
|
||||
OpSelectionMerge %144 None
|
||||
OpBranchConditional %143 %145 %144
|
||||
%145 = OpLabel
|
||||
%147 = OpCompositeExtract %float %128 1
|
||||
%148 = OpCompositeExtract %float %124 1
|
||||
%146 = OpFunctionCall %bool %aboutEqual %147 %148
|
||||
OpBranch %144
|
||||
%144 = OpLabel
|
||||
%149 = OpPhi %bool %143 %138 %146 %145
|
||||
OpStore %success %149
|
||||
OpBranch %134
|
||||
%136 = OpLabel
|
||||
%150 = OpLoad %bool %success
|
||||
OpSelectionMerge %151 None
|
||||
OpBranchConditional %150 %152 %151
|
||||
%152 = OpLabel
|
||||
%154 = OpCompositeExtract %float %128 0
|
||||
%155 = OpCompositeExtract %float %124 0
|
||||
%153 = OpFunctionCall %bool %aboutEqual %154 %155
|
||||
OpBranch %151
|
||||
%151 = OpLabel
|
||||
%156 = OpPhi %bool %150 %136 %153 %152
|
||||
OpSelectionMerge %157 None
|
||||
OpBranchConditional %156 %158 %157
|
||||
%158 = OpLabel
|
||||
%160 = OpCompositeExtract %float %128 1
|
||||
%161 = OpCompositeExtract %float %124 1
|
||||
%159 = OpFunctionCall %bool %aboutEqual %160 %161
|
||||
OpBranch %157
|
||||
%157 = OpLabel
|
||||
%162 = OpPhi %bool %156 %151 %159 %158
|
||||
OpSelectionMerge %163 None
|
||||
OpBranchConditional %162 %164 %163
|
||||
%164 = OpLabel
|
||||
%166 = OpCompositeExtract %float %128 2
|
||||
%167 = OpCompositeExtract %float %124 2
|
||||
%165 = OpFunctionCall %bool %aboutEqual %166 %167
|
||||
OpBranch %163
|
||||
%163 = OpLabel
|
||||
%168 = OpPhi %bool %162 %157 %165 %164
|
||||
OpSelectionMerge %169 None
|
||||
OpBranchConditional %168 %170 %169
|
||||
%170 = OpLabel
|
||||
%172 = OpCompositeExtract %float %128 3
|
||||
%173 = OpCompositeExtract %float %124 3
|
||||
%171 = OpFunctionCall %bool %aboutEqual %172 %173
|
||||
OpBranch %169
|
||||
%169 = OpLabel
|
||||
%174 = OpPhi %bool %168 %163 %171 %170
|
||||
OpStore %success %174
|
||||
OpBranch %134
|
||||
%134 = OpLabel
|
||||
OpBranch %87
|
||||
%87 = OpLabel
|
||||
%176 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1
|
||||
%177 = OpLoad %uint %176
|
||||
%179 = OpCompositeExtract %int %38 0
|
||||
%178 = OpBitcast %uint %179
|
||||
%180 = OpIMul %uint %177 %178
|
||||
%181 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0
|
||||
%182 = OpLoad %uint %181
|
||||
%183 = OpIAdd %uint %180 %182
|
||||
%184 = OpLoad %bool %success
|
||||
OpSelectionMerge %185 None
|
||||
OpBranchConditional %184 %186 %187
|
||||
%186 = OpLabel
|
||||
%189 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %183
|
||||
OpStore %189 %uint_1
|
||||
OpBranch %185
|
||||
%187 = OpLabel
|
||||
%190 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %183
|
||||
OpStore %190 %uint_0
|
||||
OpBranch %185
|
||||
%185 = OpLabel
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
55
test/bug/tint/913.wgsl.expected.wgsl
Normal file
55
test/bug/tint/913.wgsl.expected.wgsl
Normal file
@ -0,0 +1,55 @@
|
||||
[[block]]
|
||||
struct Uniforms {
|
||||
dstTextureFlipY : u32;
|
||||
channelCount : u32;
|
||||
srcCopyOrigin : vec2<u32>;
|
||||
dstCopyOrigin : vec2<u32>;
|
||||
copySize : vec2<u32>;
|
||||
};
|
||||
|
||||
[[block]]
|
||||
struct OutputBuf {
|
||||
result : 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 aboutEqual(value : f32, expect : f32) -> bool {
|
||||
return (abs((value - expect)) < 0.001);
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(1, 1, 1)]]
|
||||
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
||||
let srcSize : vec2<i32> = textureDimensions(src);
|
||||
let dstSize : vec2<i32> = textureDimensions(dst);
|
||||
let dstTexCoord : vec2<u32> = vec2<u32>(GlobalInvocationID.xy);
|
||||
let nonCoveredColor : vec4<f32> = vec4<f32>(0.0, 1.0, 0.0, 1.0);
|
||||
var success : bool = true;
|
||||
if (((((dstTexCoord.x < uniforms.dstCopyOrigin.x) || (dstTexCoord.y < uniforms.dstCopyOrigin.y)) || (dstTexCoord.x >= (uniforms.dstCopyOrigin.x + uniforms.copySize.x))) || (dstTexCoord.y >= (uniforms.dstCopyOrigin.y + uniforms.copySize.y)))) {
|
||||
success = (success && all((textureLoad(dst, vec2<i32>(dstTexCoord), 0) == nonCoveredColor)));
|
||||
} else {
|
||||
var srcTexCoord : vec2<u32> = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin);
|
||||
if ((uniforms.dstTextureFlipY == 1u)) {
|
||||
srcTexCoord.y = ((u32(srcSize.y) - srcTexCoord.y) - 1u);
|
||||
}
|
||||
let srcColor : vec4<f32> = textureLoad(src, vec2<i32>(srcTexCoord), 0);
|
||||
let dstColor : vec4<f32> = textureLoad(dst, vec2<i32>(dstTexCoord), 0);
|
||||
if ((uniforms.channelCount == 2u)) {
|
||||
success = ((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g));
|
||||
} else {
|
||||
success = ((((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g)) && aboutEqual(dstColor.b, srcColor.b)) && aboutEqual(dstColor.a, srcColor.a));
|
||||
}
|
||||
}
|
||||
let outputIndex : u32 = ((GlobalInvocationID.y * u32(dstSize.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
output.result[outputIndex] = 1u;
|
||||
} else {
|
||||
output.result[outputIndex] = 0u;
|
||||
}
|
||||
}
|
||||
Loading…
x
Reference in New Issue
Block a user