encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

tint/resolver: Fix use after move 

Bug: chromium:1434271
Change-Id: Ib70f9468c60fe00ca7bf9346b6ee8a71a819b59d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/127681
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
This commit is contained in:
Ben Clayton 2023-04-18 21:15:44 +00:00 committed by Dawn LUCI CQ
parent 04147fc429
commit 692846d0a0
8 changed files with 1701 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
src/tint/resolver/resolver.cc
View File

@ -139,18 +139,20 @@ bool Resolver::Resolve() {
return false; return false;
} }
// Check before std::move()'ing enabled_extensions_
const bool disable_uniformity_analysis =
enabled_extensions_.Contains(builtin::Extension::kChromiumDisableUniformityAnalysis);
// Create the semantic module. // Create the semantic module.
auto* mod = builder_->create<sem::Module>(std::move(dependencies_.ordered_globals), auto* mod = builder_->create<sem::Module>(std::move(dependencies_.ordered_globals),
std::move(enabled_extensions_)); std::move(enabled_extensions_));
ApplyDiagnosticSeverities(mod); ApplyDiagnosticSeverities(mod);
builder_->Sem().SetModule(mod); builder_->Sem().SetModule(mod);
if (result) { if (result && !disable_uniformity_analysis) {
// Run the uniformity analysis, which requires a complete semantic module. // Run the uniformity analysis, which requires a complete semantic module.
if (!enabled_extensions_.Contains(builtin::Extension::kChromiumDisableUniformityAnalysis)) { if (!AnalyzeUniformity(builder_, dependencies_)) {
if (!AnalyzeUniformity(builder_, dependencies_)) { return false;
return false;
}
} }
} }

132
test/tint/bug/chromium/1434271.wgsl Normal file
View File

@ -0,0 +1,132 @@
// Co///// * /////
enable f16;
// fn asinh(f16) -> f16
fn asinh_468a48() {
var arg_0 = f16();
var res: f16 = asinh(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
asinh_468a48();
return vec4<f32>();
}
@fragment
fn fragment_main() {
asinh_468a48();
}
@compute @workgroup_size(1)
fn rgba32uintin() {
asinh_468a48();
}
struct TestData {
dmat2atxa2 : array<atomic<i32>, 4>,
}
var<private> rand_seed : vec2<f32>;
struct RenderParams {
modelViewProjectionMatrix : mat4x4<f32>,
right : vec3<f32>,
up : vec3<f32>,
}
@binding(5) @group(0) var<uniform> render_params : RenderParams;
struct VertexInput {
@location(0)
position : vec3<f32>,
@location(1)
color : vec4<f32>,
@location(2)
quad_pos : vec2<f32>,
}
struct VertexOutput {
@builtin(position)
position : vec4<f32>,
@location(0)
color : vec4<f32>,
@location(1)
quad_pos : vec2<f32>,
}
@vertex
fn vs_main(in : VertexInput) -> VertexOutput {
var quad_pos = (mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos);
var position = (in.position - (quad_pos + 0.01));
var out : VertexOutput;
out.position = (render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0));
out.color = in.color;
out.quad_pos = in.quad_pos;
return out;
}
struct SimulationParams {
deltaTime : f32,
seed : vec4<f32>,
}
struct Particle {
position : vec3<f32>,
lifetime : f32,
color : vec4<f32>,
velocity : vec2<f32>,
}
struct Particles {
particles : array<Particle>,
}
@binding(0) @group(0) var<uniform> sim_params : SimulationParams;
@binding(1) @group(0) var<storage, read_write> data : Particles;
@binding(2) @group(0) var texture : texture_1d<f32>;
@compute @workgroup_size(64)
fn simulate(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
rand_seed = ((sim_params.seed.xy * vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw);
let idx = GlobalInvocationID.x;
var particle = data.particles[idx];
data.particles[idx] = particle;
}
struct UBO {
width : u32,
}
struct Buffer {
weights : array<f32>,
}
@binding(3) @group(0) var<uniform> ubo : UBO;
@binding(4) @group(0) var<storage, read> buf_in : Buffer;
@binding(5) @group(0) var<storage, read_write> buf_out : Buffer;
@binding(6) @group(0) var tex_in : texture_2d<f32>;
@binding(7) @group(0) var tex_out : texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(64)
fn export_level(@builtin(global_invocation_id) coord : vec3<u32>) {
if (all((coord.xy < vec2<u32>(textureDimensions(tex_out))))) {
let dst_offset = (coord.x << (coord.y * ubo.width));
let src_offset = ((coord.x - 2u) + ((coord.y >> 2u) * ubo.width));
let a = buf_in.weights[(src_offset << 0u)];
let b = buf_in.weights[(src_offset + 1u)];
let c = buf_in.weights[((src_offset + 1u) + ubo.width)];
let d = buf_in.weights[((src_offset + 1u) + ubo.width)];
let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
buf_out.weights[dst_offset] = (sum % 4.0);
let probabilities = (vec4<f32>(a, (a * b), ((a / b) + c), sum) + max(sum, 0.0));
textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
}
}

175
test/tint/bug/chromium/1434271.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,175 @@
float tint_trunc(float param_0) {
return param_0 < 0 ? ceil(param_0) : floor(param_0);
}
float16_t tint_sinh(float16_t x) {
return log((x + sqrt(((x * x) + float16_t(1.0h)))));
}
void asinh_468a48() {
float16_t arg_0 = float16_t(0.0h);
float16_t res = tint_sinh(arg_0);
}
struct tint_symbol_4 {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
asinh_468a48();
return (0.0f).xxxx;
}
tint_symbol_4 vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol_4 wrapper_result = (tint_symbol_4)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
asinh_468a48();
return;
}
[numthreads(1, 1, 1)]
void rgba32uintin() {
asinh_468a48();
return;
}
static float2 rand_seed = float2(0.0f, 0.0f);
cbuffer cbuffer_render_params : register(b5) {
uint4 render_params[6];
};
struct VertexInput {
float3 position;
float4 color;
float2 quad_pos;
};
struct VertexOutput {
float4 position;
float4 color;
float2 quad_pos;
};
struct tint_symbol_6 {
float3 position : TEXCOORD0;
float4 color : TEXCOORD1;
float2 quad_pos : TEXCOORD2;
};
struct tint_symbol_7 {
float4 color : TEXCOORD0;
float2 quad_pos : TEXCOORD1;
float4 position : SV_Position;
};
float4x4 render_params_load_1(uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
const uint scalar_offset_1 = ((offset + 16u)) / 4;
const uint scalar_offset_2 = ((offset + 32u)) / 4;
const uint scalar_offset_3 = ((offset + 48u)) / 4;
return float4x4(asfloat(render_params[scalar_offset / 4]), asfloat(render_params[scalar_offset_1 / 4]), asfloat(render_params[scalar_offset_2 / 4]), asfloat(render_params[scalar_offset_3 / 4]));
}
VertexOutput vs_main_inner(VertexInput tint_symbol) {
float3 quad_pos = mul(tint_symbol.quad_pos, float2x3(asfloat(render_params[4].xyz), asfloat(render_params[5].xyz)));
float3 position = (tint_symbol.position - (quad_pos + 0.00999999977648258209f));
VertexOutput tint_symbol_1 = (VertexOutput)0;
tint_symbol_1.position = mul(float4(position, 1.0f), render_params_load_1(0u));
tint_symbol_1.color = tint_symbol.color;
tint_symbol_1.quad_pos = tint_symbol.quad_pos;
return tint_symbol_1;
}
tint_symbol_7 vs_main(tint_symbol_6 tint_symbol_5) {
const VertexInput tint_symbol_12 = {tint_symbol_5.position, tint_symbol_5.color, tint_symbol_5.quad_pos};
const VertexOutput inner_result_1 = vs_main_inner(tint_symbol_12);
tint_symbol_7 wrapper_result_1 = (tint_symbol_7)0;
wrapper_result_1.position = inner_result_1.position;
wrapper_result_1.color = inner_result_1.color;
wrapper_result_1.quad_pos = inner_result_1.quad_pos;
return wrapper_result_1;
}
struct Particle {
float3 position;
float lifetime;
float4 color;
float2 velocity;
};
cbuffer cbuffer_sim_params : register(b0) {
uint4 sim_params[2];
};
RWByteAddressBuffer data : register(u1);
Texture1D<float4> tint_symbol_2 : register(t2);
struct tint_symbol_9 {
uint3 GlobalInvocationID : SV_DispatchThreadID;
};
Particle data_load(uint offset) {
const Particle tint_symbol_13 = {asfloat(data.Load3((offset + 0u))), asfloat(data.Load((offset + 12u))), asfloat(data.Load4((offset + 16u))), asfloat(data.Load2((offset + 32u)))};
return tint_symbol_13;
}
void data_store(uint offset, Particle value) {
data.Store3((offset + 0u), asuint(value.position));
data.Store((offset + 12u), asuint(value.lifetime));
data.Store4((offset + 16u), asuint(value.color));
data.Store2((offset + 32u), asuint(value.velocity));
}
void simulate_inner(uint3 GlobalInvocationID) {
rand_seed = ((asfloat(sim_params[1]).xy * float2(GlobalInvocationID.xy)) * asfloat(sim_params[1]).zw);
const uint idx = GlobalInvocationID.x;
Particle particle = data_load((48u * idx));
data_store((48u * idx), particle);
}
[numthreads(64, 1, 1)]
void simulate(tint_symbol_9 tint_symbol_8) {
simulate_inner(tint_symbol_8.GlobalInvocationID);
return;
}
cbuffer cbuffer_ubo : register(b3) {
uint4 ubo[1];
};
ByteAddressBuffer buf_in : register(t4);
RWByteAddressBuffer buf_out : register(u5);
Texture2D<float4> tex_in : register(t6);
RWTexture2D<float4> tex_out : register(u7);
float tint_float_mod(float lhs, float rhs) {
return (lhs - (tint_trunc((lhs / rhs)) * rhs));
}
struct tint_symbol_11 {
uint3 coord : SV_DispatchThreadID;
};
void export_level_inner(uint3 coord) {
uint2 tint_tmp;
tex_out.GetDimensions(tint_tmp.x, tint_tmp.y);
if (all((coord.xy < uint2(tint_tmp)))) {
const uint dst_offset = (coord.x << ((coord.y * ubo[0].x) & 31u));
const uint src_offset = ((coord.x - 2u) + ((coord.y >> 2u) * ubo[0].x));
const float a = asfloat(buf_in.Load((4u * (src_offset << 0u))));
const float b = asfloat(buf_in.Load((4u * (src_offset + 1u))));
const float c = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
const float d = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
const float sum = dot(float4(a, b, c, d), (1.0f).xxxx);
buf_out.Store((4u * dst_offset), asuint(tint_float_mod(sum, 4.0f)));
const float4 probabilities = (float4(a, (a * b), ((a / b) + c), sum) + max(sum, 0.0f));
tex_out[int2(coord.xy)] = probabilities;
}
}
[numthreads(64, 1, 1)]
void export_level(tint_symbol_11 tint_symbol_10) {
export_level_inner(tint_symbol_10.coord);
return;
}

177
test/tint/bug/chromium/1434271.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,177 @@
SKIP: f16 not supported
float tint_trunc(float param_0) {
return param_0 < 0 ? ceil(param_0) : floor(param_0);
}
float16_t tint_sinh(float16_t x) {
return log((x + sqrt(((x * x) + float16_t(1.0h)))));
}
void asinh_468a48() {
float16_t arg_0 = float16_t(0.0h);
float16_t res = tint_sinh(arg_0);
}
struct tint_symbol_4 {
float4 value : SV_Position;
};
float4 vertex_main_inner() {
asinh_468a48();
return (0.0f).xxxx;
}
tint_symbol_4 vertex_main() {
const float4 inner_result = vertex_main_inner();
tint_symbol_4 wrapper_result = (tint_symbol_4)0;
wrapper_result.value = inner_result;
return wrapper_result;
}
void fragment_main() {
asinh_468a48();
return;
}
[numthreads(1, 1, 1)]
void rgba32uintin() {
asinh_468a48();
return;
}
static float2 rand_seed = float2(0.0f, 0.0f);
cbuffer cbuffer_render_params : register(b5) {
uint4 render_params[6];
};
struct VertexInput {
float3 position;
float4 color;
float2 quad_pos;
};
struct VertexOutput {
float4 position;
float4 color;
float2 quad_pos;
};
struct tint_symbol_6 {
float3 position : TEXCOORD0;
float4 color : TEXCOORD1;
float2 quad_pos : TEXCOORD2;
};
struct tint_symbol_7 {
float4 color : TEXCOORD0;
float2 quad_pos : TEXCOORD1;
float4 position : SV_Position;
};
float4x4 render_params_load_1(uint offset) {
const uint scalar_offset = ((offset + 0u)) / 4;
const uint scalar_offset_1 = ((offset + 16u)) / 4;
const uint scalar_offset_2 = ((offset + 32u)) / 4;
const uint scalar_offset_3 = ((offset + 48u)) / 4;
return float4x4(asfloat(render_params[scalar_offset / 4]), asfloat(render_params[scalar_offset_1 / 4]), asfloat(render_params[scalar_offset_2 / 4]), asfloat(render_params[scalar_offset_3 / 4]));
}
VertexOutput vs_main_inner(VertexInput tint_symbol) {
float3 quad_pos = mul(tint_symbol.quad_pos, float2x3(asfloat(render_params[4].xyz), asfloat(render_params[5].xyz)));
float3 position = (tint_symbol.position - (quad_pos + 0.00999999977648258209f));
VertexOutput tint_symbol_1 = (VertexOutput)0;
tint_symbol_1.position = mul(float4(position, 1.0f), render_params_load_1(0u));
tint_symbol_1.color = tint_symbol.color;
tint_symbol_1.quad_pos = tint_symbol.quad_pos;
return tint_symbol_1;
}
tint_symbol_7 vs_main(tint_symbol_6 tint_symbol_5) {
const VertexInput tint_symbol_12 = {tint_symbol_5.position, tint_symbol_5.color, tint_symbol_5.quad_pos};
const VertexOutput inner_result_1 = vs_main_inner(tint_symbol_12);
tint_symbol_7 wrapper_result_1 = (tint_symbol_7)0;
wrapper_result_1.position = inner_result_1.position;
wrapper_result_1.color = inner_result_1.color;
wrapper_result_1.quad_pos = inner_result_1.quad_pos;
return wrapper_result_1;
}
struct Particle {
float3 position;
float lifetime;
float4 color;
float2 velocity;
};
cbuffer cbuffer_sim_params : register(b0) {
uint4 sim_params[2];
};
RWByteAddressBuffer data : register(u1);
Texture1D<float4> tint_symbol_2 : register(t2);
struct tint_symbol_9 {
uint3 GlobalInvocationID : SV_DispatchThreadID;
};
Particle data_load(uint offset) {
const Particle tint_symbol_13 = {asfloat(data.Load3((offset + 0u))), asfloat(data.Load((offset + 12u))), asfloat(data.Load4((offset + 16u))), asfloat(data.Load2((offset + 32u)))};
return tint_symbol_13;
}
void data_store(uint offset, Particle value) {
data.Store3((offset + 0u), asuint(value.position));
data.Store((offset + 12u), asuint(value.lifetime));
data.Store4((offset + 16u), asuint(value.color));
data.Store2((offset + 32u), asuint(value.velocity));
}
void simulate_inner(uint3 GlobalInvocationID) {
rand_seed = ((asfloat(sim_params[1]).xy * float2(GlobalInvocationID.xy)) * asfloat(sim_params[1]).zw);
const uint idx = GlobalInvocationID.x;
Particle particle = data_load((48u * idx));
data_store((48u * idx), particle);
}
[numthreads(64, 1, 1)]
void simulate(tint_symbol_9 tint_symbol_8) {
simulate_inner(tint_symbol_8.GlobalInvocationID);
return;
}
cbuffer cbuffer_ubo : register(b3) {
uint4 ubo[1];
};
ByteAddressBuffer buf_in : register(t4);
RWByteAddressBuffer buf_out : register(u5);
Texture2D<float4> tex_in : register(t6);
RWTexture2D<float4> tex_out : register(u7);
float tint_float_mod(float lhs, float rhs) {
return (lhs - (tint_trunc((lhs / rhs)) * rhs));
}
struct tint_symbol_11 {
uint3 coord : SV_DispatchThreadID;
};
void export_level_inner(uint3 coord) {
uint2 tint_tmp;
tex_out.GetDimensions(tint_tmp.x, tint_tmp.y);
if (all((coord.xy < uint2(tint_tmp)))) {
const uint dst_offset = (coord.x << ((coord.y * ubo[0].x) & 31u));
const uint src_offset = ((coord.x - 2u) + ((coord.y >> 2u) * ubo[0].x));
const float a = asfloat(buf_in.Load((4u * (src_offset << 0u))));
const float b = asfloat(buf_in.Load((4u * (src_offset + 1u))));
const float c = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
const float d = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
const float sum = dot(float4(a, b, c, d), (1.0f).xxxx);
buf_out.Store((4u * dst_offset), asuint(tint_float_mod(sum, 4.0f)));
const float4 probabilities = (float4(a, (a * b), ((a / b) + c), sum) + max(sum, 0.0f));
tex_out[int2(coord.xy)] = probabilities;
}
}
[numthreads(64, 1, 1)]
void export_level(tint_symbol_11 tint_symbol_10) {
export_level_inner(tint_symbol_10.coord);
return;
}

399
test/tint/bug/chromium/1434271.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,399 @@
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
void asinh_468a48() {
float16_t arg_0 = 0.0hf;
float16_t res = asinh(arg_0);
}
struct TestData {
int dmat2atxa2[4];
};
struct RenderParams {
mat4 modelViewProjectionMatrix;
vec3 right;
vec3 up;
};
struct VertexInput {
vec3 position;
vec4 color;
vec2 quad_pos;
};
struct VertexOutput {
vec4 position;
vec4 color;
vec2 quad_pos;
};
struct SimulationParams {
float deltaTime;
vec4 seed;
};
struct Particle {
vec3 position;
float lifetime;
vec4 color;
vec2 velocity;
};
struct UBO {
uint width;
};
vec4 vertex_main() {
asinh_468a48();
return vec4(0.0f);
}
void main() {
gl_PointSize = 1.0;
vec4 inner_result = vertex_main();
gl_Position = inner_result;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
precision highp float;
void asinh_468a48() {
float16_t arg_0 = 0.0hf;
float16_t res = asinh(arg_0);
}
struct TestData {
int dmat2atxa2[4];
};
struct RenderParams {
mat4 modelViewProjectionMatrix;
vec3 right;
vec3 up;
};
struct VertexInput {
vec3 position;
vec4 color;
vec2 quad_pos;
};
struct VertexOutput {
vec4 position;
vec4 color;
vec2 quad_pos;
};
struct SimulationParams {
float deltaTime;
vec4 seed;
};
struct Particle {
vec3 position;
float lifetime;
vec4 color;
vec2 velocity;
};
struct UBO {
uint width;
};
void fragment_main() {
asinh_468a48();
}
void main() {
fragment_main();
return;
}
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
void asinh_468a48() {
float16_t arg_0 = 0.0hf;
float16_t res = asinh(arg_0);
}
struct TestData {
int dmat2atxa2[4];
};
struct RenderParams {
mat4 modelViewProjectionMatrix;
vec3 right;
vec3 up;
};
struct VertexInput {
vec3 position;
vec4 color;
vec2 quad_pos;
};
struct VertexOutput {
vec4 position;
vec4 color;
vec2 quad_pos;
};
struct SimulationParams {
float deltaTime;
vec4 seed;
};
struct Particle {
vec3 position;
float lifetime;
vec4 color;
vec2 velocity;
};
struct UBO {
uint width;
};
void rgba32uintin() {
asinh_468a48();
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
rgba32uintin();
return;
}
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
layout(location = 0) in vec3 position_1;
layout(location = 1) in vec4 color_1;
layout(location = 2) in vec2 quad_pos_1;
layout(location = 0) out vec4 color_2;
layout(location = 1) out vec2 quad_pos_2;
struct TestData {
int dmat2atxa2[4];
};
struct RenderParams {
mat4 modelViewProjectionMatrix;
vec3 right;
uint pad;
vec3 up;
uint pad_1;
};
layout(binding = 5, std140) uniform render_params_block_ubo {
RenderParams inner;
} render_params;
struct VertexInput {
vec3 position;
vec4 color;
vec2 quad_pos;
};
struct VertexOutput {
vec4 position;
vec4 color;
vec2 quad_pos;
};
struct SimulationParams {
float deltaTime;
vec4 seed;
};
struct Particle {
vec3 position;
float lifetime;
vec4 color;
vec2 velocity;
};
struct UBO {
uint width;
};
VertexOutput vs_main(VertexInput tint_symbol) {
vec3 quad_pos = (mat2x3(render_params.inner.right, render_params.inner.up) * tint_symbol.quad_pos);
vec3 position = (tint_symbol.position - (quad_pos + 0.00999999977648258209f));
VertexOutput tint_symbol_1 = VertexOutput(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec4(0.0f, 0.0f, 0.0f, 0.0f), vec2(0.0f, 0.0f));
tint_symbol_1.position = (render_params.inner.modelViewProjectionMatrix * vec4(position, 1.0f));
tint_symbol_1.color = tint_symbol.color;
tint_symbol_1.quad_pos = tint_symbol.quad_pos;
return tint_symbol_1;
}
void main() {
gl_PointSize = 1.0;
VertexInput tint_symbol_2 = VertexInput(position_1, color_1, quad_pos_1);
VertexOutput inner_result = vs_main(tint_symbol_2);
gl_Position = inner_result.position;
color_2 = inner_result.color;
quad_pos_2 = inner_result.quad_pos;
gl_Position.y = -(gl_Position.y);
gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
return;
}
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
struct TestData {
int dmat2atxa2[4];
};
vec2 rand_seed = vec2(0.0f, 0.0f);
struct RenderParams {
mat4 modelViewProjectionMatrix;
vec3 right;
vec3 up;
};
struct VertexInput {
vec3 position;
vec4 color;
vec2 quad_pos;
};
struct VertexOutput {
vec4 position;
vec4 color;
vec2 quad_pos;
};
struct SimulationParams {
float deltaTime;
uint pad;
uint pad_1;
uint pad_2;
vec4 seed;
};
struct Particle {
vec3 position;
float lifetime;
vec4 color;
vec2 velocity;
uint pad_3;
uint pad_4;
};
layout(binding = 0, std140) uniform sim_params_block_ubo {
SimulationParams inner;
} sim_params;
layout(binding = 1, std430) buffer Particles_ssbo {
Particle particles[];
} data;
struct UBO {
uint width;
};
void assign_and_preserve_padding_data_particles_X(uint dest[1], Particle value) {
data.particles[dest[0]].position = value.position;
data.particles[dest[0]].lifetime = value.lifetime;
data.particles[dest[0]].color = value.color;
data.particles[dest[0]].velocity = value.velocity;
}
void simulate(uvec3 GlobalInvocationID) {
rand_seed = ((sim_params.inner.seed.xy * vec2(GlobalInvocationID.xy)) * sim_params.inner.seed.zw);
uint idx = GlobalInvocationID.x;
Particle particle = data.particles[idx];
uint tint_symbol[1] = uint[1](idx);
assign_and_preserve_padding_data_particles_X(tint_symbol, particle);
}
layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
void main() {
simulate(gl_GlobalInvocationID);
return;
}
#version 310 es
#extension GL_AMD_gpu_shader_half_float : require
float tint_float_modulo(float lhs, float rhs) {
return (lhs - rhs * trunc(lhs / rhs));
}
struct TestData {
int dmat2atxa2[4];
};
struct RenderParams {
mat4 modelViewProjectionMatrix;
vec3 right;
vec3 up;
};
struct VertexInput {
vec3 position;
vec4 color;
vec2 quad_pos;
};
struct VertexOutput {
vec4 position;
vec4 color;
vec2 quad_pos;
};
struct SimulationParams {
float deltaTime;
vec4 seed;
};
struct Particle {
vec3 position;
float lifetime;
vec4 color;
vec2 velocity;
};
struct UBO {
uint width;
uint pad;
uint pad_1;
uint pad_2;
};
layout(binding = 3, std140) uniform ubo_block_ubo {
UBO inner;
} ubo;
layout(binding = 4, std430) buffer Buffer_ssbo {
float weights[];
} buf_in;
layout(binding = 5, std430) buffer Buffer_ssbo_1 {
float weights[];
} buf_out;
layout(rgba8) uniform highp writeonly image2D tex_out;
void export_level(uvec3 coord) {
if (all(lessThan(coord.xy, uvec2(uvec2(imageSize(tex_out)))))) {
uint dst_offset = (coord.x << ((coord.y * ubo.inner.width) & 31u));
uint src_offset = ((coord.x - 2u) + ((coord.y >> 2u) * ubo.inner.width));
float a = buf_in.weights[(src_offset << 0u)];
float b = buf_in.weights[(src_offset + 1u)];
float c = buf_in.weights[((src_offset + 1u) + ubo.inner.width)];
float d = buf_in.weights[((src_offset + 1u) + ubo.inner.width)];
float sum = dot(vec4(a, b, c, d), vec4(1.0f));
buf_out.weights[dst_offset] = tint_float_modulo(sum, 4.0f);
vec4 probabilities = (vec4(a, (a * b), ((a / b) + c), sum) + max(sum, 0.0f));
imageStore(tex_out, ivec2(coord.xy), probabilities);
}
}
layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
void main() {
export_level(gl_GlobalInvocationID);
return;
}

199
test/tint/bug/chromium/1434271.wgsl.expected.msl Normal file
View File

@ -0,0 +1,199 @@
#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
struct tint_private_vars_struct {
float2 rand_seed;
};
struct RenderParams_tint_packed_vec3 {
/* 0x0000 */ float4x4 modelViewProjectionMatrix;
/* 0x0040 */ packed_float3 right;
/* 0x004c */ tint_array<int8_t, 4> tint_pad;
/* 0x0050 */ packed_float3 up;
/* 0x005c */ tint_array<int8_t, 4> tint_pad_1;
};
struct Particle_tint_packed_vec3 {
/* 0x0000 */ packed_float3 position;
/* 0x000c */ float lifetime;
/* 0x0010 */ float4 color;
/* 0x0020 */ float2 velocity;
/* 0x0028 */ tint_array<int8_t, 8> tint_pad_2;
};
struct Particles_tint_packed_vec3 {
/* 0x0000 */ tint_array<Particle_tint_packed_vec3, 1> particles;
};
struct Particle {
float3 position;
float lifetime;
float4 color;
float2 velocity;
};
Particle tint_unpack_vec3_in_composite(Particle_tint_packed_vec3 in) {
Particle result = {};
result.position = float3(in.position);
result.lifetime = in.lifetime;
result.color = in.color;
result.velocity = in.velocity;
return result;
}
void asinh_468a48() {
half arg_0 = 0.0h;
half res = asinh(arg_0);
}
struct tint_symbol_1 {
float4 value [[position]];
};
float4 vertex_main_inner() {
asinh_468a48();
return float4(0.0f);
}
vertex tint_symbol_1 vertex_main() {
float4 const inner_result = vertex_main_inner();
tint_symbol_1 wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
fragment void fragment_main() {
asinh_468a48();
return;
}
kernel void rgba32uintin() {
asinh_468a48();
return;
}
struct TestData {
tint_array<atomic_int, 4> dmat2atxa2;
};
struct RenderParams {
float4x4 modelViewProjectionMatrix;
float3 right;
float3 up;
};
struct VertexInput {
float3 position;
float4 color;
float2 quad_pos;
};
struct VertexOutput {
float4 position;
float4 color;
float2 quad_pos;
};
struct tint_symbol_3 {
float3 position [[attribute(0)]];
float4 color [[attribute(1)]];
float2 quad_pos [[attribute(2)]];
};
struct tint_symbol_4 {
float4 color [[user(locn0)]];
float2 quad_pos [[user(locn1)]];
float4 position [[position]];
};
VertexOutput vs_main_inner(VertexInput in, const constant RenderParams_tint_packed_vec3* const tint_symbol_6) {
float3 quad_pos = (float2x3(float3((*(tint_symbol_6)).right), float3((*(tint_symbol_6)).up)) * in.quad_pos);
float3 position = (in.position - (quad_pos + 0.00999999977648258209f));
VertexOutput out = {};
out.position = ((*(tint_symbol_6)).modelViewProjectionMatrix * float4(position, 1.0f));
out.color = in.color;
out.quad_pos = in.quad_pos;
return out;
}
vertex tint_symbol_4 vs_main(const constant RenderParams_tint_packed_vec3* tint_symbol_7 [[buffer(0)]], tint_symbol_3 tint_symbol_2 [[stage_in]]) {
VertexInput const tint_symbol_5 = {.position=tint_symbol_2.position, .color=tint_symbol_2.color, .quad_pos=tint_symbol_2.quad_pos};
VertexOutput const inner_result_1 = vs_main_inner(tint_symbol_5, tint_symbol_7);
tint_symbol_4 wrapper_result_1 = {};
wrapper_result_1.position = inner_result_1.position;
wrapper_result_1.color = inner_result_1.color;
wrapper_result_1.quad_pos = inner_result_1.quad_pos;
return wrapper_result_1;
}
struct SimulationParams {
/* 0x0000 */ float deltaTime;
/* 0x0004 */ tint_array<int8_t, 12> tint_pad_3;
/* 0x0010 */ float4 seed;
};
struct Particles {
tint_array<Particle, 1> particles;
};
void assign_and_preserve_padding(device Particle_tint_packed_vec3* const dest, Particle value) {
(*(dest)).position = packed_float3(value.position);
(*(dest)).lifetime = value.lifetime;
(*(dest)).color = value.color;
(*(dest)).velocity = value.velocity;
}
void simulate_inner(uint3 GlobalInvocationID, thread tint_private_vars_struct* const tint_private_vars, const constant SimulationParams* const tint_symbol_8, device Particles_tint_packed_vec3* const tint_symbol_9) {
(*(tint_private_vars)).rand_seed = (((*(tint_symbol_8)).seed.xy * float2(GlobalInvocationID.xy)) * (*(tint_symbol_8)).seed.zw);
uint const idx = GlobalInvocationID[0];
Particle particle = tint_unpack_vec3_in_composite((*(tint_symbol_9)).particles[idx]);
assign_and_preserve_padding(&((*(tint_symbol_9)).particles[idx]), particle);
}
kernel void simulate(const constant SimulationParams* tint_symbol_10 [[buffer(1)]], device Particles_tint_packed_vec3* tint_symbol_11 [[buffer(2)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
thread tint_private_vars_struct tint_private_vars = {};
simulate_inner(GlobalInvocationID, &(tint_private_vars), tint_symbol_10, tint_symbol_11);
return;
}
struct UBO {
/* 0x0000 */ uint width;
};
struct Buffer {
/* 0x0000 */ tint_array<float, 1> weights;
};
void export_level_inner(uint3 coord, texture2d<float, access::write> tint_symbol_12, const constant UBO* const tint_symbol_13, const device Buffer* const tint_symbol_14, device Buffer* const tint_symbol_15) {
if (all((coord.xy < uint2(uint2(tint_symbol_12.get_width(), tint_symbol_12.get_height()))))) {
uint const dst_offset = (coord[0] << ((coord[1] * (*(tint_symbol_13)).width) & 31u));
uint const src_offset = ((coord[0] - 2u) + ((coord[1] >> 2u) * (*(tint_symbol_13)).width));
float const a = (*(tint_symbol_14)).weights[(src_offset << 0u)];
float const b = (*(tint_symbol_14)).weights[(src_offset + 1u)];
float const c = (*(tint_symbol_14)).weights[((src_offset + 1u) + (*(tint_symbol_13)).width)];
float const d = (*(tint_symbol_14)).weights[((src_offset + 1u) + (*(tint_symbol_13)).width)];
float const sum = dot(float4(a, b, c, d), float4(1.0f));
(*(tint_symbol_15)).weights[dst_offset] = fmod(sum, 4.0f);
float4 const probabilities = (float4(a, (a * b), ((a / b) + c), sum) + fmax(sum, 0.0f));
tint_symbol_12.write(probabilities, uint2(int2(coord.xy)));
}
}
kernel void export_level(texture2d<float, access::write> tint_symbol_16 [[texture(0)]], const constant UBO* tint_symbol_17 [[buffer(3)]], const device Buffer* tint_symbol_18 [[buffer(4)]], device Buffer* tint_symbol_19 [[buffer(0)]], uint3 coord [[thread_position_in_grid]]) {
export_level_inner(coord, tint_symbol_16, tint_symbol_17, tint_symbol_18, tint_symbol_19);
return;
}

483
test/tint/bug/chromium/1434271.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,483 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 274
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
OpCapability Sampled1D
OpCapability ImageQuery
%71 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value_1 %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
OpEntryPoint GLCompute %rgba32uintin "rgba32uintin"
OpEntryPoint Vertex %vs_main "vs_main" %position_1 %color_1 %quad_pos_1 %position_2 %color_2 %quad_pos_2 %vertex_point_size_1
OpEntryPoint GLCompute %simulate "simulate" %GlobalInvocationID_1
OpEntryPoint GLCompute %export_level "export_level" %coord_1
OpExecutionMode %fragment_main OriginUpperLeft
OpExecutionMode %rgba32uintin LocalSize 1 1 1
OpExecutionMode %simulate LocalSize 64 1 1
OpExecutionMode %export_level LocalSize 64 1 1
OpName %value_1 "value_1"
OpName %vertex_point_size "vertex_point_size"
OpName %position_1 "position_1"
OpName %color_1 "color_1"
OpName %quad_pos_1 "quad_pos_1"
OpName %position_2 "position_2"
OpName %color_2 "color_2"
OpName %quad_pos_2 "quad_pos_2"
OpName %vertex_point_size_1 "vertex_point_size_1"
OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
OpName %coord_1 "coord_1"
OpName %rand_seed "rand_seed"
OpName %render_params_block "render_params_block"
OpMemberName %render_params_block 0 "inner"
OpName %RenderParams "RenderParams"
OpMemberName %RenderParams 0 "modelViewProjectionMatrix"
OpMemberName %RenderParams 1 "right"
OpMemberName %RenderParams 2 "up"
OpName %render_params "render_params"
OpName %sim_params_block "sim_params_block"
OpMemberName %sim_params_block 0 "inner"
OpName %SimulationParams "SimulationParams"
OpMemberName %SimulationParams 0 "deltaTime"
OpMemberName %SimulationParams 1 "seed"
OpName %sim_params "sim_params"
OpName %Particles "Particles"
OpMemberName %Particles 0 "particles"
OpName %Particle "Particle"
OpMemberName %Particle 0 "position"
OpMemberName %Particle 1 "lifetime"
OpMemberName %Particle 2 "color"
OpMemberName %Particle 3 "velocity"
OpName %data "data"
OpName %texture "texture"
OpName %ubo_block "ubo_block"
OpMemberName %ubo_block 0 "inner"
OpName %UBO "UBO"
OpMemberName %UBO 0 "width"
OpName %ubo "ubo"
OpName %Buffer "Buffer"
OpMemberName %Buffer 0 "weights"
OpName %buf_in "buf_in"
OpName %buf_out "buf_out"
OpName %tex_in "tex_in"
OpName %tex_out "tex_out"
OpName %asinh_468a48 "asinh_468a48"
OpName %arg_0 "arg_0"
OpName %res "res"
OpName %vertex_main_inner "vertex_main_inner"
OpName %vertex_main "vertex_main"
OpName %fragment_main "fragment_main"
OpName %rgba32uintin "rgba32uintin"
OpName %VertexOutput "VertexOutput"
OpMemberName %VertexOutput 0 "position"
OpMemberName %VertexOutput 1 "color"
OpMemberName %VertexOutput 2 "quad_pos"
OpName %VertexInput "VertexInput"
OpMemberName %VertexInput 0 "position"
OpMemberName %VertexInput 1 "color"
OpMemberName %VertexInput 2 "quad_pos"
OpName %vs_main_inner "vs_main_inner"
OpName %in "in"
OpName %quad_pos "quad_pos"
OpName %position "position"
OpName %out "out"
OpName %vs_main "vs_main"
OpName %assign_and_preserve_padding_data_particles_X "assign_and_preserve_padding_data_particles_X"
OpName %dest "dest"
OpName %value "value"
OpName %simulate_inner "simulate_inner"
OpName %GlobalInvocationID "GlobalInvocationID"
OpName %particle "particle"
OpName %simulate "simulate"
OpName %export_level_inner "export_level_inner"
OpName %coord "coord"
OpName %export_level "export_level"
OpDecorate %value_1 BuiltIn Position
OpDecorate %vertex_point_size BuiltIn PointSize
OpDecorate %position_1 Location 0
OpDecorate %color_1 Location 1
OpDecorate %quad_pos_1 Location 2
OpDecorate %position_2 BuiltIn Position
OpDecorate %color_2 Location 0
OpDecorate %quad_pos_2 Location 1
OpDecorate %vertex_point_size_1 BuiltIn PointSize
OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
OpDecorate %coord_1 BuiltIn GlobalInvocationId
OpDecorate %render_params_block Block
OpMemberDecorate %render_params_block 0 Offset 0
OpMemberDecorate %RenderParams 0 Offset 0
OpMemberDecorate %RenderParams 0 ColMajor
OpMemberDecorate %RenderParams 0 MatrixStride 16
OpMemberDecorate %RenderParams 1 Offset 64
OpMemberDecorate %RenderParams 2 Offset 80
OpDecorate %render_params NonWritable
OpDecorate %render_params Binding 5
OpDecorate %render_params DescriptorSet 0
OpDecorate %sim_params_block Block
OpMemberDecorate %sim_params_block 0 Offset 0
OpMemberDecorate %SimulationParams 0 Offset 0
OpMemberDecorate %SimulationParams 1 Offset 16
OpDecorate %sim_params NonWritable
OpDecorate %sim_params Binding 0
OpDecorate %sim_params DescriptorSet 0
OpDecorate %Particles Block
OpMemberDecorate %Particles 0 Offset 0
OpMemberDecorate %Particle 0 Offset 0
OpMemberDecorate %Particle 1 Offset 12
OpMemberDecorate %Particle 2 Offset 16
OpMemberDecorate %Particle 3 Offset 32
OpDecorate %_runtimearr_Particle ArrayStride 48
OpDecorate %data Binding 1
OpDecorate %data DescriptorSet 0
OpDecorate %texture Binding 2
OpDecorate %texture DescriptorSet 0
OpDecorate %ubo_block Block
OpMemberDecorate %ubo_block 0 Offset 0
OpMemberDecorate %UBO 0 Offset 0
OpDecorate %ubo NonWritable
OpDecorate %ubo Binding 3
OpDecorate %ubo DescriptorSet 0
OpDecorate %Buffer Block
OpMemberDecorate %Buffer 0 Offset 0
OpDecorate %_runtimearr_float ArrayStride 4
OpDecorate %buf_in NonWritable
OpDecorate %buf_in Binding 4
OpDecorate %buf_in DescriptorSet 0
OpDecorate %buf_out Binding 5
OpDecorate %buf_out DescriptorSet 0
OpDecorate %tex_in Binding 6
OpDecorate %tex_in DescriptorSet 0
OpDecorate %tex_out NonReadable
OpDecorate %tex_out Binding 7
OpDecorate %tex_out DescriptorSet 0
OpMemberDecorate %VertexOutput 0 Offset 0
OpMemberDecorate %VertexOutput 1 Offset 16
OpMemberDecorate %VertexOutput 2 Offset 32
OpMemberDecorate %VertexInput 0 Offset 0
OpMemberDecorate %VertexInput 1 Offset 16
OpMemberDecorate %VertexInput 2 Offset 32
OpDecorate %_arr_uint_uint_1 ArrayStride 4
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value_1 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_float = OpTypePointer Output %float
%8 = OpConstantNull %float
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%v3float = OpTypeVector %float 3
%_ptr_Input_v3float = OpTypePointer Input %v3float
%position_1 = OpVariable %_ptr_Input_v3float Input
%_ptr_Input_v4float = OpTypePointer Input %v4float
%color_1 = OpVariable %_ptr_Input_v4float Input
%v2float = OpTypeVector %float 2
%_ptr_Input_v2float = OpTypePointer Input %v2float
%quad_pos_1 = OpVariable %_ptr_Input_v2float Input
%position_2 = OpVariable %_ptr_Output_v4float Output %5
%color_2 = OpVariable %_ptr_Output_v4float Output %5
%_ptr_Output_v2float = OpTypePointer Output %v2float
%21 = OpConstantNull %v2float
%quad_pos_2 = OpVariable %_ptr_Output_v2float Output %21
%vertex_point_size_1 = OpVariable %_ptr_Output_float Output %8
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
%coord_1 = OpVariable %_ptr_Input_v3uint Input
%_ptr_Private_v2float = OpTypePointer Private %v2float
%rand_seed = OpVariable %_ptr_Private_v2float Private %21
%mat4v4float = OpTypeMatrix %v4float 4
%RenderParams = OpTypeStruct %mat4v4float %v3float %v3float
%render_params_block = OpTypeStruct %RenderParams
%_ptr_Uniform_render_params_block = OpTypePointer Uniform %render_params_block
%render_params = OpVariable %_ptr_Uniform_render_params_block Uniform
%SimulationParams = OpTypeStruct %float %v4float
%sim_params_block = OpTypeStruct %SimulationParams
%_ptr_Uniform_sim_params_block = OpTypePointer Uniform %sim_params_block
%sim_params = OpVariable %_ptr_Uniform_sim_params_block Uniform
%Particle = OpTypeStruct %v3float %float %v4float %v2float
%_runtimearr_Particle = OpTypeRuntimeArray %Particle
%Particles = OpTypeStruct %_runtimearr_Particle
%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
%data = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
%46 = OpTypeImage %float 1D 0 0 0 1 Unknown
%_ptr_UniformConstant_46 = OpTypePointer UniformConstant %46
%texture = OpVariable %_ptr_UniformConstant_46 UniformConstant
%UBO = OpTypeStruct %uint
%ubo_block = OpTypeStruct %UBO
%_ptr_Uniform_ubo_block = OpTypePointer Uniform %ubo_block
%ubo = OpVariable %_ptr_Uniform_ubo_block Uniform
%_runtimearr_float = OpTypeRuntimeArray %float
%Buffer = OpTypeStruct %_runtimearr_float
%_ptr_StorageBuffer_Buffer = OpTypePointer StorageBuffer %Buffer
%buf_in = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
%buf_out = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
%58 = OpTypeImage %float 2D 0 0 0 1 Unknown
%_ptr_UniformConstant_58 = OpTypePointer UniformConstant %58
%tex_in = OpVariable %_ptr_UniformConstant_58 UniformConstant
%61 = OpTypeImage %float 2D 0 0 0 2 Rgba8
%_ptr_UniformConstant_61 = OpTypePointer UniformConstant %61
%tex_out = OpVariable %_ptr_UniformConstant_61 UniformConstant
%void = OpTypeVoid
%62 = OpTypeFunction %void
%half = OpTypeFloat 16
%67 = OpConstantNull %half
%_ptr_Function_half = OpTypePointer Function %half
%74 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%VertexOutput = OpTypeStruct %v4float %v4float %v2float
%VertexInput = OpTypeStruct %v3float %v4float %v2float
%88 = OpTypeFunction %VertexOutput %VertexInput
%mat2v3float = OpTypeMatrix %v3float 2
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%uint_2 = OpConstant %uint 2
%_ptr_Function_v3float = OpTypePointer Function %v3float
%108 = OpConstantNull %v3float
%float_0_00999999978 = OpConstant %float 0.00999999978
%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
%119 = OpConstantNull %VertexOutput
%_ptr_Function_v4float = OpTypePointer Function %v4float
%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
%_ptr_Function_v2float = OpTypePointer Function %v2float
%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
%147 = OpTypeFunction %void %_arr_uint_uint_1 %Particle
%int = OpTypeInt 32 1
%154 = OpConstantNull %int
%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
%uint_3 = OpConstant %uint 3
%_ptr_StorageBuffer_v2float = OpTypePointer StorageBuffer %v2float
%172 = OpTypeFunction %void %v3uint
%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
%v2uint = OpTypeVector %uint 2
%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle
%_ptr_Function_Particle = OpTypePointer Function %Particle
%194 = OpConstantNull %Particle
%bool = OpTypeBool
%v2bool = OpTypeVector %bool 2
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%uint_31 = OpConstant %uint 31
%232 = OpConstantNull %uint
%253 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
%float_4 = OpConstant %float 4
%v2int = OpTypeVector %int 2
%asinh_468a48 = OpFunction %void None %62
%65 = OpLabel
%arg_0 = OpVariable %_ptr_Function_half Function %67
%res = OpVariable %_ptr_Function_half Function %67
OpStore %arg_0 %67
%72 = OpLoad %half %arg_0
%70 = OpExtInst %half %71 Asinh %72
OpStore %res %70
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %74
%76 = OpLabel
%77 = OpFunctionCall %void %asinh_468a48
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %62
%79 = OpLabel
%80 = OpFunctionCall %v4float %vertex_main_inner
OpStore %value_1 %80
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %62
%83 = OpLabel
%84 = OpFunctionCall %void %asinh_468a48
OpReturn
OpFunctionEnd
%rgba32uintin = OpFunction %void None %62
%86 = OpLabel
%87 = OpFunctionCall %void %asinh_468a48
OpReturn
OpFunctionEnd
%vs_main_inner = OpFunction %VertexOutput None %88
%in = OpFunctionParameter %VertexInput
%93 = OpLabel
%quad_pos = OpVariable %_ptr_Function_v3float Function %108
%113 = OpVariable %_ptr_Function_v3float Function %108
%position = OpVariable %_ptr_Function_v3float Function %108
%out = OpVariable %_ptr_Function_VertexOutput Function %119
%98 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_0 %uint_1
%99 = OpLoad %v3float %98
%101 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_0 %uint_2
%102 = OpLoad %v3float %101
%103 = OpCompositeConstruct %mat2v3float %99 %102
%104 = OpCompositeExtract %v2float %in 2
%105 = OpMatrixTimesVector %v3float %103 %104
OpStore %quad_pos %105
%109 = OpCompositeExtract %v3float %in 0
%110 = OpLoad %v3float %quad_pos
%114 = OpCompositeConstruct %v3float %float_0_00999999978 %float_0_00999999978 %float_0_00999999978
%112 = OpFAdd %v3float %110 %114
%115 = OpFSub %v3float %109 %112
OpStore %position %115
%121 = OpAccessChain %_ptr_Function_v4float %out %uint_0
%123 = OpAccessChain %_ptr_Uniform_mat4v4float %render_params %uint_0 %uint_0
%124 = OpLoad %mat4v4float %123
%125 = OpLoad %v3float %position
%126 = OpCompositeExtract %float %125 0
%127 = OpCompositeExtract %float %125 1
%128 = OpCompositeExtract %float %125 2
%129 = OpCompositeConstruct %v4float %126 %127 %128 %float_1
%130 = OpMatrixTimesVector %v4float %124 %129
OpStore %121 %130
%131 = OpAccessChain %_ptr_Function_v4float %out %uint_1
%132 = OpCompositeExtract %v4float %in 1
OpStore %131 %132
%134 = OpAccessChain %_ptr_Function_v2float %out %uint_2
%135 = OpCompositeExtract %v2float %in 2
OpStore %134 %135
%136 = OpLoad %VertexOutput %out
OpReturnValue %136
OpFunctionEnd
%vs_main = OpFunction %void None %62
%138 = OpLabel
%140 = OpLoad %v3float %position_1
%141 = OpLoad %v4float %color_1
%142 = OpLoad %v2float %quad_pos_1
%143 = OpCompositeConstruct %VertexInput %140 %141 %142
%139 = OpFunctionCall %VertexOutput %vs_main_inner %143
%144 = OpCompositeExtract %v4float %139 0
OpStore %position_2 %144
%145 = OpCompositeExtract %v4float %139 1
OpStore %color_2 %145
%146 = OpCompositeExtract %v2float %139 2
OpStore %quad_pos_2 %146
OpStore %vertex_point_size_1 %float_1
OpReturn
OpFunctionEnd
%assign_and_preserve_padding_data_particles_X = OpFunction %void None %147
%dest = OpFunctionParameter %_arr_uint_uint_1
%value = OpFunctionParameter %Particle
%152 = OpLabel
%155 = OpCompositeExtract %uint %dest 0
%157 = OpAccessChain %_ptr_StorageBuffer_v3float %data %uint_0 %155 %uint_0
%158 = OpCompositeExtract %v3float %value 0
OpStore %157 %158
%159 = OpCompositeExtract %uint %dest 0
%161 = OpAccessChain %_ptr_StorageBuffer_float %data %uint_0 %159 %uint_1
%162 = OpCompositeExtract %float %value 1
OpStore %161 %162
%163 = OpCompositeExtract %uint %dest 0
%165 = OpAccessChain %_ptr_StorageBuffer_v4float %data %uint_0 %163 %uint_2
%166 = OpCompositeExtract %v4float %value 2
OpStore %165 %166
%167 = OpCompositeExtract %uint %dest 0
%170 = OpAccessChain %_ptr_StorageBuffer_v2float %data %uint_0 %167 %uint_3
%171 = OpCompositeExtract %v2float %value 3
OpStore %170 %171
OpReturn
OpFunctionEnd
%simulate_inner = OpFunction %void None %172
%GlobalInvocationID = OpFunctionParameter %v3uint
%175 = OpLabel
%particle = OpVariable %_ptr_Function_Particle Function %194
%177 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_0 %uint_1
%178 = OpLoad %v4float %177
%179 = OpVectorShuffle %v2float %178 %178 0 1
%182 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
%180 = OpConvertUToF %v2float %182
%183 = OpFMul %v2float %179 %180
%184 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_0 %uint_1
%185 = OpLoad %v4float %184
%186 = OpVectorShuffle %v2float %185 %185 2 3
%187 = OpFMul %v2float %183 %186
OpStore %rand_seed %187
%188 = OpCompositeExtract %uint %GlobalInvocationID 0
%190 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
%191 = OpLoad %Particle %190
OpStore %particle %191
%196 = OpCompositeConstruct %_arr_uint_uint_1 %188
%197 = OpLoad %Particle %particle
%195 = OpFunctionCall %void %assign_and_preserve_padding_data_particles_X %196 %197
OpReturn
OpFunctionEnd
%simulate = OpFunction %void None %62
%199 = OpLabel
%201 = OpLoad %v3uint %GlobalInvocationID_1
%200 = OpFunctionCall %void %simulate_inner %201
OpReturn
OpFunctionEnd
%export_level_inner = OpFunction %void None %172
%coord = OpFunctionParameter %v3uint
%204 = OpLabel
%263 = OpVariable %_ptr_Function_v4float Function %5
%207 = OpVectorShuffle %v2uint %coord %coord 0 1
%210 = OpLoad %61 %tex_out
%209 = OpImageQuerySize %v2uint %210
%211 = OpULessThan %v2bool %207 %209
%205 = OpAll %bool %211
OpSelectionMerge %213 None
OpBranchConditional %205 %214 %213
%214 = OpLabel
%215 = OpCompositeExtract %uint %coord 0
%216 = OpCompositeExtract %uint %coord 1
%218 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0 %uint_0
%219 = OpLoad %uint %218
%220 = OpIMul %uint %216 %219
%222 = OpBitwiseAnd %uint %220 %uint_31
%223 = OpShiftLeftLogical %uint %215 %222
%224 = OpCompositeExtract %uint %coord 0
%225 = OpISub %uint %224 %uint_2
%226 = OpCompositeExtract %uint %coord 1
%227 = OpShiftRightLogical %uint %226 %uint_2
%228 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0 %uint_0
%229 = OpLoad %uint %228
%230 = OpIMul %uint %227 %229
%231 = OpIAdd %uint %225 %230
%233 = OpShiftLeftLogical %uint %231 %232
%234 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %233
%235 = OpLoad %float %234
%236 = OpIAdd %uint %231 %uint_1
%237 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %236
%238 = OpLoad %float %237
%239 = OpIAdd %uint %231 %uint_1
%240 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0 %uint_0
%241 = OpLoad %uint %240
%242 = OpIAdd %uint %239 %241
%243 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %242
%244 = OpLoad %float %243
%245 = OpIAdd %uint %231 %uint_1
%246 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0 %uint_0
%247 = OpLoad %uint %246
%248 = OpIAdd %uint %245 %247
%249 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %248
%250 = OpLoad %float %249
%252 = OpCompositeConstruct %v4float %235 %238 %244 %250
%251 = OpDot %float %252 %253
%254 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %223
%256 = OpFRem %float %251 %float_4
OpStore %254 %256
%257 = OpFMul %float %235 %238
%258 = OpFDiv %float %235 %238
%259 = OpFAdd %float %258 %244
%260 = OpCompositeConstruct %v4float %235 %257 %259 %251
%261 = OpExtInst %float %71 NMax %251 %8
%264 = OpCompositeConstruct %v4float %261 %261 %261 %261
%262 = OpFAdd %v4float %260 %264
%266 = OpLoad %61 %tex_out
%269 = OpVectorShuffle %v2uint %coord %coord 0 1
%267 = OpBitcast %v2int %269
OpImageWrite %266 %267 %262
OpBranch %213
%213 = OpLabel
OpReturn
OpFunctionEnd
%export_level = OpFunction %void None %62
%271 = OpLabel
%273 = OpLoad %v3uint %coord_1
%272 = OpFunctionCall %void %export_level_inner %273
OpReturn
OpFunctionEnd

129
test/tint/bug/chromium/1434271.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,129 @@
enable f16;
fn asinh_468a48() {
var arg_0 = f16();
var res : f16 = asinh(arg_0);
}
@vertex
fn vertex_main() -> @builtin(position) vec4<f32> {
asinh_468a48();
return vec4<f32>();
}
@fragment
fn fragment_main() {
asinh_468a48();
}
@compute @workgroup_size(1)
fn rgba32uintin() {
asinh_468a48();
}
struct TestData {
dmat2atxa2 : array<atomic<i32>, 4>,
}
var<private> rand_seed : vec2<f32>;
struct RenderParams {
modelViewProjectionMatrix : mat4x4<f32>,
right : vec3<f32>,
up : vec3<f32>,
}
@binding(5) @group(0) var<uniform> render_params : RenderParams;
struct VertexInput {
@location(0)
position : vec3<f32>,
@location(1)
color : vec4<f32>,
@location(2)
quad_pos : vec2<f32>,
}
struct VertexOutput {
@builtin(position)
position : vec4<f32>,
@location(0)
color : vec4<f32>,
@location(1)
quad_pos : vec2<f32>,
}
@vertex
fn vs_main(in : VertexInput) -> VertexOutput {
var quad_pos = (mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos);
var position = (in.position - (quad_pos + 0.01000000000000000021));
var out : VertexOutput;
out.position = (render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0));
out.color = in.color;
out.quad_pos = in.quad_pos;
return out;
}
struct SimulationParams {
deltaTime : f32,
seed : vec4<f32>,
}
struct Particle {
position : vec3<f32>,
lifetime : f32,
color : vec4<f32>,
velocity : vec2<f32>,
}
struct Particles {
particles : array<Particle>,
}
@binding(0) @group(0) var<uniform> sim_params : SimulationParams;
@binding(1) @group(0) var<storage, read_write> data : Particles;
@binding(2) @group(0) var texture : texture_1d<f32>;
@compute @workgroup_size(64)
fn simulate(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
rand_seed = ((sim_params.seed.xy * vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw);
let idx = GlobalInvocationID.x;
var particle = data.particles[idx];
data.particles[idx] = particle;
}
struct UBO {
width : u32,
}
struct Buffer {
weights : array<f32>,
}
@binding(3) @group(0) var<uniform> ubo : UBO;
@binding(4) @group(0) var<storage, read> buf_in : Buffer;
@binding(5) @group(0) var<storage, read_write> buf_out : Buffer;
@binding(6) @group(0) var tex_in : texture_2d<f32>;
@binding(7) @group(0) var tex_out : texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(64)
fn export_level(@builtin(global_invocation_id) coord : vec3<u32>) {
if (all((coord.xy < vec2<u32>(textureDimensions(tex_out))))) {
let dst_offset = (coord.x << (coord.y * ubo.width));
let src_offset = ((coord.x - 2u) + ((coord.y >> 2u) * ubo.width));
let a = buf_in.weights[(src_offset << 0u)];
let b = buf_in.weights[(src_offset + 1u)];
let c = buf_in.weights[((src_offset + 1u) + ubo.width)];
let d = buf_in.weights[((src_offset + 1u) + ubo.width)];
let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
buf_out.weights[dst_offset] = (sum % 4.0);
let probabilities = (vec4<f32>(a, (a * b), ((a / b) + c), sum) + max(sum, 0.0));
textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
}
}