writer/hlsl: UnwrapRef() on atomic types
Fixed: tint:1113
Change-Id: I9aa255f5b308cc4d53b0ea40407cc398096a502c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/61780
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index c36db0d..12449fb 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -993,7 +993,7 @@
case Op::kAtomicStore: {
// HLSL does not have an InterlockedStore, so we emulate it with
// InterlockedExchange and discard the returned value
- auto* value_ty = TypeOf(expr->params()[2]);
+ auto* value_ty = TypeOf(expr->params()[2])->UnwrapRef();
auto name = UniqueIdentifier("atomicStore");
{
auto fn = line(&buf);
@@ -1024,7 +1024,7 @@
return name;
}
case Op::kAtomicCompareExchangeWeak: {
- auto* value_ty = TypeOf(expr->params()[2]);
+ auto* value_ty = TypeOf(expr->params()[2])->UnwrapRef();
auto name = UniqueIdentifier("atomicCompareExchangeWeak");
{
@@ -1167,7 +1167,7 @@
// InterlockedExchange and discard the returned value
{ // T result = 0;
auto pre = line();
- auto* value_ty = intrinsic->Parameters()[1]->Type();
+ auto* value_ty = intrinsic->Parameters()[1]->Type()->UnwrapRef();
if (!EmitTypeAndName(pre, value_ty, ast::StorageClass::kNone,
ast::Access::kUndefined, result)) {
return false;
diff --git a/test/bug/tint/1113.wgsl b/test/bug/tint/1113.wgsl
new file mode 100644
index 0000000..dba7d01
--- /dev/null
+++ b/test/bug/tint/1113.wgsl
@@ -0,0 +1,184 @@
+
+[[block]] struct Uniforms {
+ numTriangles : u32;
+ gridSize : u32;
+ pad1 : u32;
+ pad2 : u32;
+ bbMin : vec3<f32>; // offset(16)
+ bbMax : vec3<f32>; // offset(32)
+};
+
+[[block]] struct Dbg {
+ offsetCounter : atomic<u32>;
+ pad0 : u32;
+ pad1 : u32;
+ pad2 : u32;
+ value0 : u32;
+ value1 : u32;
+ value2 : u32;
+ value3 : u32;
+ value_f32_0 : f32;
+ value_f32_1 : f32;
+ value_f32_2 : f32;
+ value_f32_3 : f32;
+};
+
+[[block]] struct F32s { values : [[stride(4)]] array<f32>; };
+[[block]] struct U32s { values : [[stride(4)]] array<u32>; };
+[[block]] struct I32s { values : [[stride(4)]] array<i32>; };
+[[block]] struct AU32s { values : [[stride(4)]] array<atomic<u32>>; };
+[[block]] struct AI32s { values : [[stride(4)]] array<atomic<i32>>; };
+
+// IN
+[[binding(0), group(0)]] var<uniform> uniforms : Uniforms;
+[[binding(10), group(0)]] var<storage, read_write> indices : U32s;
+[[binding(11), group(0)]] var<storage, read_write> positions : F32s;
+
+// OUT
+[[binding(20), group(0)]] var<storage, read_write> counters : AU32s;
+[[binding(21), group(0)]] var<storage, read_write> LUT : AI32s;
+
+// DEBUG
+[[binding(50), group(0)]] var<storage, read_write> dbg : Dbg;
+
+fn toVoxelPos(position : vec3<f32>) -> vec3<f32>{
+
+ var bbMin = vec3<f32>(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z);
+ var bbMax = vec3<f32>(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z);
+ var bbSize = bbMax - bbMin;
+ var cubeSize = max(max(bbSize.x, bbSize.y), bbSize.z);
+ var gridSize = f32(uniforms.gridSize);
+
+ var gx = gridSize * (position.x - uniforms.bbMin.x) / cubeSize;
+ var gy = gridSize * (position.y - uniforms.bbMin.y) / cubeSize;
+ var gz = gridSize * (position.z - uniforms.bbMin.z) / cubeSize;
+
+ return vec3<f32>(gx, gy, gz);
+}
+
+fn toIndex1D(gridSize : u32, voxelPos : vec3<f32>) -> u32{
+
+ var icoord = vec3<u32>(voxelPos);
+
+ return icoord.x
+ + gridSize * icoord.y
+ + gridSize * gridSize * icoord.z;
+}
+
+fn toIndex3D(gridSize : u32, index : u32) -> vec3<u32>{
+ var z = index / (gridSize * gridSize);
+ var y = (index - gridSize * gridSize * z) / gridSize;
+ var x = index % gridSize;
+
+ return vec3<u32>(x, y, z);
+}
+
+fn loadPosition(vertexIndex : u32) -> vec3<f32> {
+
+ var position = vec3<f32>(
+ positions.values[3u * vertexIndex + 0u],
+ positions.values[3u * vertexIndex + 1u],
+ positions.values[3u * vertexIndex + 2u],
+ );
+
+ return position;
+};
+
+fn doIgnore(){
+
+ var g42 = uniforms.numTriangles;
+ var kj6 = dbg.value1;
+ var b53 = atomicLoad(&counters.values[0]);
+ var rwg = indices.values[0];
+ var rb5 = positions.values[0];
+ var g55 = atomicLoad(&LUT.values[0]);
+
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_count([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+
+ var triangleIndex = GlobalInvocationID.x;
+
+ if(triangleIndex >= uniforms.numTriangles){
+ return;
+ }
+
+ doIgnore();
+
+ var i0 = indices.values[3u * triangleIndex + 0u];
+ var i1 = indices.values[3u * triangleIndex + 1u];
+ var i2 = indices.values[3u * triangleIndex + 2u];
+
+ var p0 = loadPosition(i0);
+ var p1 = loadPosition(i1);
+ var p2 = loadPosition(i2);
+ var center = (p0 + p1 + p2) / 3.0;
+
+ var voxelPos = toVoxelPos(center);
+ var voxelIndex = toIndex1D(uniforms.gridSize, voxelPos);
+
+ var acefg = atomicAdd(&counters.values[voxelIndex], 1u);
+
+ if(triangleIndex == 0u){
+ dbg.value0 = uniforms.gridSize;
+ // dbg.value1 = voxelPos.y;
+ // dbg.value2 = voxelPos.z;
+
+ dbg.value_f32_0 = center.x;
+ dbg.value_f32_1 = center.y;
+ dbg.value_f32_2 = center.z;
+ }
+
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_create_lut([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+
+ var voxelIndex = GlobalInvocationID.x;
+
+ doIgnore();
+
+ var maxVoxels = uniforms.gridSize * uniforms.gridSize * uniforms.gridSize;
+ if(voxelIndex >= maxVoxels){
+ return;
+ }
+
+ var numTriangles = atomicLoad(&counters.values[voxelIndex]);
+
+ var offset = -1;
+ if(numTriangles > 0u){
+ offset = i32(atomicAdd(&dbg.offsetCounter, numTriangles));
+ }
+
+ atomicStore(&LUT.values[voxelIndex], offset);
+
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_sort_triangles([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+
+ var triangleIndex = GlobalInvocationID.x;
+
+ doIgnore();
+
+ if(triangleIndex >= uniforms.numTriangles){
+ return;
+ }
+
+ var i0 = indices.values[3u * triangleIndex + 0u];
+ var i1 = indices.values[3u * triangleIndex + 1u];
+ var i2 = indices.values[3u * triangleIndex + 2u];
+
+ var p0 = loadPosition(i0);
+ var p1 = loadPosition(i1);
+ var p2 = loadPosition(i2);
+ var center = (p0 + p1 + p2) / 3.0;
+
+ var voxelPos = toVoxelPos(center);
+ var voxelIndex = toIndex1D(uniforms.gridSize, voxelPos);
+
+ // var triangleOffset = LUT.values[voxelIndex];
+ var triangleOffset = atomicAdd(&LUT.values[voxelIndex], 1);
+
+}
diff --git a/test/bug/tint/1113.wgsl.expected.hlsl b/test/bug/tint/1113.wgsl.expected.hlsl
new file mode 100644
index 0000000..e630a4c
--- /dev/null
+++ b/test/bug/tint/1113.wgsl.expected.hlsl
@@ -0,0 +1,162 @@
+uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
+ uint value = 0;
+ buffer.InterlockedOr(offset, 0, value);
+ return value;
+}
+
+int atomicLoad_2(RWByteAddressBuffer buffer, uint offset) {
+ int value = 0;
+ buffer.InterlockedOr(offset, 0, value);
+ return value;
+}
+
+uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+ uint original_value = 0;
+ buffer.InterlockedAdd(offset, value, original_value);
+ return original_value;
+}
+
+void atomicStore_1(RWByteAddressBuffer buffer, uint offset, int value) {
+ int ignored;
+ buffer.InterlockedExchange(offset, value, ignored);
+}
+
+int atomicAdd_2(RWByteAddressBuffer buffer, uint offset, int value) {
+ int original_value = 0;
+ buffer.InterlockedAdd(offset, value, original_value);
+ return original_value;
+}
+
+cbuffer cbuffer_uniforms : register(b0, space0) {
+ uint4 uniforms[3];
+};
+RWByteAddressBuffer indices : register(u10, space0);
+RWByteAddressBuffer positions : register(u11, space0);
+RWByteAddressBuffer counters : register(u20, space0);
+RWByteAddressBuffer LUT : register(u21, space0);
+RWByteAddressBuffer dbg : register(u50, space0);
+
+float3 toVoxelPos(float3 position) {
+ float3 bbMin = float3(asfloat(uniforms[1].x), asfloat(uniforms[1].y), asfloat(uniforms[1].z));
+ float3 bbMax = float3(asfloat(uniforms[2].x), asfloat(uniforms[2].y), asfloat(uniforms[2].z));
+ float3 bbSize = (bbMax - bbMin);
+ float cubeSize = max(max(bbSize.x, bbSize.y), bbSize.z);
+ float gridSize = float(uniforms[0].y);
+ float gx = ((gridSize * (position.x - asfloat(uniforms[1].x))) / cubeSize);
+ float gy = ((gridSize * (position.y - asfloat(uniforms[1].y))) / cubeSize);
+ float gz = ((gridSize * (position.z - asfloat(uniforms[1].z))) / cubeSize);
+ return float3(gx, gy, gz);
+}
+
+uint toIndex1D(uint gridSize, float3 voxelPos) {
+ uint3 icoord = uint3(voxelPos);
+ return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
+}
+
+uint3 toIndex3D(uint gridSize, uint index) {
+ uint z_1 = (index / (gridSize * gridSize));
+ uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / gridSize);
+ uint x_1 = (index % gridSize);
+ return uint3(x_1, y_1, z_1);
+}
+
+float3 loadPosition(uint vertexIndex) {
+ float3 position = float3(asfloat(positions.Load((4u * ((3u * vertexIndex) + 0u)))), asfloat(positions.Load((4u * ((3u * vertexIndex) + 1u)))), asfloat(positions.Load((4u * ((3u * vertexIndex) + 2u)))));
+ return position;
+}
+
+void doIgnore() {
+ uint g42 = uniforms[0].x;
+ uint kj6 = dbg.Load(20u);
+ uint b53 = atomicLoad_1(counters, (4u * uint(0)));
+ uint rwg = indices.Load((4u * uint(0)));
+ float rb5 = asfloat(positions.Load((4u * uint(0))));
+ int g55 = atomicLoad_2(LUT, (4u * uint(0)));
+}
+
+struct tint_symbol_1 {
+ uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void main_count_inner(uint3 GlobalInvocationID) {
+ uint triangleIndex = GlobalInvocationID.x;
+ if ((triangleIndex >= uniforms[0].x)) {
+ return;
+ }
+ doIgnore();
+ uint i0 = indices.Load((4u * ((3u * triangleIndex) + 0u)));
+ uint i1 = indices.Load((4u * ((3u * triangleIndex) + 1u)));
+ uint i2 = indices.Load((4u * ((3u * triangleIndex) + 2u)));
+ float3 p0 = loadPosition(i0);
+ float3 p1 = loadPosition(i1);
+ float3 p2 = loadPosition(i2);
+ float3 center = (((p0 + p1) + p2) / 3.0f);
+ float3 voxelPos = toVoxelPos(center);
+ uint voxelIndex = toIndex1D(uniforms[0].y, voxelPos);
+ uint acefg = atomicAdd_1(counters, (4u * voxelIndex), 1u);
+ if ((triangleIndex == 0u)) {
+ dbg.Store(16u, asuint(uniforms[0].y));
+ dbg.Store(32u, asuint(center.x));
+ dbg.Store(36u, asuint(center.y));
+ dbg.Store(40u, asuint(center.z));
+ }
+}
+
+[numthreads(128, 1, 1)]
+void main_count(tint_symbol_1 tint_symbol) {
+ main_count_inner(tint_symbol.GlobalInvocationID);
+ return;
+}
+
+struct tint_symbol_3 {
+ uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void main_create_lut_inner(uint3 GlobalInvocationID) {
+ uint voxelIndex = GlobalInvocationID.x;
+ doIgnore();
+ uint maxVoxels = ((uniforms[0].y * uniforms[0].y) * uniforms[0].y);
+ if ((voxelIndex >= maxVoxels)) {
+ return;
+ }
+ uint numTriangles = atomicLoad_1(counters, (4u * voxelIndex));
+ int offset = -1;
+ if ((numTriangles > 0u)) {
+ offset = int(atomicAdd_1(dbg, 0u, numTriangles));
+ }
+ atomicStore_1(LUT, (4u * voxelIndex), offset);
+}
+
+[numthreads(128, 1, 1)]
+void main_create_lut(tint_symbol_3 tint_symbol_2) {
+ main_create_lut_inner(tint_symbol_2.GlobalInvocationID);
+ return;
+}
+
+struct tint_symbol_5 {
+ uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+void main_sort_triangles_inner(uint3 GlobalInvocationID) {
+ uint triangleIndex = GlobalInvocationID.x;
+ doIgnore();
+ if ((triangleIndex >= uniforms[0].x)) {
+ return;
+ }
+ uint i0 = indices.Load((4u * ((3u * triangleIndex) + 0u)));
+ uint i1 = indices.Load((4u * ((3u * triangleIndex) + 1u)));
+ uint i2 = indices.Load((4u * ((3u * triangleIndex) + 2u)));
+ float3 p0 = loadPosition(i0);
+ float3 p1 = loadPosition(i1);
+ float3 p2 = loadPosition(i2);
+ float3 center = (((p0 + p1) + p2) / 3.0f);
+ float3 voxelPos = toVoxelPos(center);
+ uint voxelIndex = toIndex1D(uniforms[0].y, voxelPos);
+ int triangleOffset = atomicAdd_2(LUT, (4u * voxelIndex), 1);
+}
+
+[numthreads(128, 1, 1)]
+void main_sort_triangles(tint_symbol_5 tint_symbol_4) {
+ main_sort_triangles_inner(tint_symbol_4.GlobalInvocationID);
+ return;
+}
diff --git a/test/bug/tint/1113.wgsl.expected.msl b/test/bug/tint/1113.wgsl.expected.msl
new file mode 100644
index 0000000..2cfcdc0
--- /dev/null
+++ b/test/bug/tint/1113.wgsl.expected.msl
@@ -0,0 +1,201 @@
+SKIP: FAILED
+
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+ /* 0x0000 */ uint numTriangles;
+ /* 0x0004 */ uint gridSize;
+ /* 0x0008 */ uint pad1;
+ /* 0x000c */ uint pad2;
+ /* 0x0010 */ packed_float3 bbMin;
+ /* 0x001c */ int8_t tint_pad[4];
+ /* 0x0020 */ packed_float3 bbMax;
+ /* 0x002c */ int8_t tint_pad_1[4];
+};
+struct Dbg {
+ /* 0x0000 */ atomic_uint offsetCounter;
+ /* 0x0004 */ uint pad0;
+ /* 0x0008 */ uint pad1;
+ /* 0x000c */ uint pad2;
+ /* 0x0010 */ uint value0;
+ /* 0x0014 */ uint value1;
+ /* 0x0018 */ uint value2;
+ /* 0x001c */ uint value3;
+ /* 0x0020 */ float value_f32_0;
+ /* 0x0024 */ float value_f32_1;
+ /* 0x0028 */ float value_f32_2;
+ /* 0x002c */ float value_f32_3;
+};
+struct F32s {
+ /* 0x0000 */ float values[1];
+};
+struct U32s {
+ /* 0x0000 */ uint values[1];
+};
+struct I32s {
+ int values[1];
+};
+struct AU32s {
+ /* 0x0000 */ atomic_uint values[1];
+};
+struct AI32s {
+ /* 0x0000 */ atomic_int values[1];
+};
+
+float3 toVoxelPos(constant Uniforms& uniforms, float3 position) {
+ float3 bbMin = float3(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z);
+ float3 bbMax = float3(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z);
+ float3 bbSize = (bbMax - bbMin);
+ float cubeSize = fmax(fmax(bbSize.x, bbSize.y), bbSize.z);
+ float gridSize = float(uniforms.gridSize);
+ float gx = ((gridSize * (position.x - uniforms.bbMin.x)) / cubeSize);
+ float gy = ((gridSize * (position.y - uniforms.bbMin.y)) / cubeSize);
+ float gz = ((gridSize * (position.z - uniforms.bbMin.z)) / cubeSize);
+ return float3(gx, gy, gz);
+}
+
+uint toIndex1D(uint gridSize, float3 voxelPos) {
+ uint3 icoord = uint3(voxelPos);
+ return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
+}
+
+uint3 toIndex3D(uint gridSize, uint index) {
+ uint z_1 = (index / (gridSize * gridSize));
+ uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / gridSize);
+ uint x_1 = (index % gridSize);
+ return uint3(x_1, y_1, z_1);
+}
+
+float3 loadPosition(device F32s& positions, uint vertexIndex) {
+ float3 position = float3(positions.values[((3u * vertexIndex) + 0u)], positions.values[((3u * vertexIndex) + 1u)], positions.values[((3u * vertexIndex) + 2u)]);
+ return position;
+}
+
+void doIgnore(constant Uniforms& uniforms, device Dbg& dbg, device AU32s& counters, device U32s& indices, device F32s& positions, device AI32s& LUT) {
+ uint g42 = uniforms.numTriangles;
+ uint kj6 = dbg.value1;
+ uint b53 = atomic_load_explicit(&(counters.values[0]), memory_order_relaxed);
+ uint rwg = indices.values[0];
+ float rb5 = positions.values[0];
+ int g55 = atomic_load_explicit(&(LUT.values[0]), memory_order_relaxed);
+}
+
+void main_count_inner(constant Uniforms& uniforms, device Dbg& dbg, device AU32s& counters, device U32s& indices, device F32s& positions, device AI32s& LUT, uint3 GlobalInvocationID) {
+ uint triangleIndex = GlobalInvocationID.x;
+ if ((triangleIndex >= uniforms.numTriangles)) {
+ return;
+ }
+ doIgnore(uniforms, dbg, counters, indices, positions, LUT);
+ uint i0 = indices.values[((3u * triangleIndex) + 0u)];
+ uint i1 = indices.values[((3u * triangleIndex) + 1u)];
+ uint i2 = indices.values[((3u * triangleIndex) + 2u)];
+ float3 p0 = loadPosition(positions, i0);
+ float3 p1 = loadPosition(positions, i1);
+ float3 p2 = loadPosition(positions, i2);
+ float3 center = (((p0 + p1) + p2) / 3.0f);
+ float3 voxelPos = toVoxelPos(uniforms, center);
+ uint voxelIndex = toIndex1D(uniforms.gridSize, voxelPos);
+ uint acefg = atomic_fetch_add_explicit(&(counters.values[voxelIndex]), 1u, memory_order_relaxed);
+ if ((triangleIndex == 0u)) {
+ dbg.value0 = uniforms.gridSize;
+ dbg.value_f32_0 = center.x;
+ dbg.value_f32_1 = center.y;
+ dbg.value_f32_2 = center.z;
+ }
+}
+
+kernel void main_count(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], device Dbg& dbg [[buffer(50)]], device AU32s& counters [[buffer(20)]], device U32s& indices [[buffer(10)]], device F32s& positions [[buffer(11)]], device AI32s& LUT [[buffer(21)]]) {
+ main_count_inner(uniforms, dbg, counters, indices, positions, LUT, GlobalInvocationID);
+ return;
+}
+
+void main_create_lut_inner(constant Uniforms& uniforms, device Dbg& dbg, device AU32s& counters, device U32s& indices, device F32s& positions, device AI32s& LUT, uint3 GlobalInvocationID) {
+ uint voxelIndex = GlobalInvocationID.x;
+ doIgnore(uniforms, dbg, counters, indices, positions, LUT);
+ uint maxVoxels = ((uniforms.gridSize * uniforms.gridSize) * uniforms.gridSize);
+ if ((voxelIndex >= maxVoxels)) {
+ return;
+ }
+ uint numTriangles = atomic_load_explicit(&(counters.values[voxelIndex]), memory_order_relaxed);
+ int offset = -1;
+ if ((numTriangles > 0u)) {
+ offset = int(atomic_fetch_add_explicit(&(dbg.offsetCounter), numTriangles, memory_order_relaxed));
+ }
+ atomic_store_explicit(&(LUT.values[voxelIndex]), offset, memory_order_relaxed);
+}
+
+kernel void main_create_lut(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], device Dbg& dbg [[buffer(50)]], device AU32s& counters [[buffer(20)]], device U32s& indices [[buffer(10)]], device F32s& positions [[buffer(11)]], device AI32s& LUT [[buffer(21)]]) {
+ main_create_lut_inner(uniforms, dbg, counters, indices, positions, LUT, GlobalInvocationID);
+ return;
+}
+
+void main_sort_triangles_inner(constant Uniforms& uniforms, device Dbg& dbg, device AU32s& counters, device U32s& indices, device F32s& positions, device AI32s& LUT, uint3 GlobalInvocationID) {
+ uint triangleIndex = GlobalInvocationID.x;
+ doIgnore(uniforms, dbg, counters, indices, positions, LUT);
+ if ((triangleIndex >= uniforms.numTriangles)) {
+ return;
+ }
+ uint i0 = indices.values[((3u * triangleIndex) + 0u)];
+ uint i1 = indices.values[((3u * triangleIndex) + 1u)];
+ uint i2 = indices.values[((3u * triangleIndex) + 2u)];
+ float3 p0 = loadPosition(positions, i0);
+ float3 p1 = loadPosition(positions, i1);
+ float3 p2 = loadPosition(positions, i2);
+ float3 center = (((p0 + p1) + p2) / 3.0f);
+ float3 voxelPos = toVoxelPos(uniforms, center);
+ uint voxelIndex = toIndex1D(uniforms.gridSize, voxelPos);
+ int triangleOffset = atomic_fetch_add_explicit(&(LUT.values[voxelIndex]), 1, memory_order_relaxed);
+}
+
+kernel void main_sort_triangles(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], device Dbg& dbg [[buffer(50)]], device AU32s& counters [[buffer(20)]], device U32s& indices [[buffer(10)]], device F32s& positions [[buffer(11)]], device AI32s& LUT [[buffer(21)]]) {
+ main_sort_triangles_inner(uniforms, dbg, counters, indices, positions, LUT, GlobalInvocationID);
+ return;
+}
+
+Compilation failed:
+
+program_source:75:8: warning: unused variable 'kj6'
+ uint kj6 = dbg.value1;
+ ^
+program_source:78:9: warning: unused variable 'rb5'
+ float rb5 = positions.values[0];
+ ^
+program_source:79:7: warning: unused variable 'g55'
+ int g55 = atomic_load_explicit(&(LUT.values[0]), memory_order_relaxed);
+ ^
+program_source:77:8: warning: unused variable 'rwg'
+ uint rwg = indices.values[0];
+ ^
+program_source:74:8: warning: unused variable 'g42'
+ uint g42 = uniforms.numTriangles;
+ ^
+program_source:76:8: warning: unused variable 'b53'
+ uint b53 = atomic_load_explicit(&(counters.values[0]), memory_order_relaxed);
+ ^
+program_source:98:22: warning: equality comparison with extraneous parentheses
+ if ((triangleIndex == 0u)) {
+ ~~~~~~~~~~~~~~^~~~~
+program_source:98:22: note: remove extraneous parentheses around the comparison to silence this warning
+ if ((triangleIndex == 0u)) {
+ ~ ^ ~
+program_source:98:22: note: use '=' to turn this equality comparison into an assignment
+ if ((triangleIndex == 0u)) {
+ ^~
+ =
+program_source:97:8: warning: unused variable 'acefg'
+ uint acefg = atomic_fetch_add_explicit(&(counters.values[voxelIndex]), 1u, memory_order_relaxed);
+ ^
+program_source:106:146: error: 'buffer' attribute parameter is out of bounds: must be between 0 and 30
+kernel void main_count(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], device Dbg& dbg [[buffer(50)]], device AU32s& counters [[buffer(20)]], device U32s& indices [[buffer(10)]], device F32s& positions [[buffer(11)]], device AI32s& LUT [[buffer(21)]]) {
+ ^
+program_source:126:151: error: 'buffer' attribute parameter is out of bounds: must be between 0 and 30
+kernel void main_create_lut(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], device Dbg& dbg [[buffer(50)]], device AU32s& counters [[buffer(20)]], device U32s& indices [[buffer(10)]], device F32s& positions [[buffer(11)]], device AI32s& LUT [[buffer(21)]]) {
+ ^
+program_source:146:7: warning: unused variable 'triangleOffset'
+ int triangleOffset = atomic_fetch_add_explicit(&(LUT.values[voxelIndex]), 1, memory_order_relaxed);
+ ^
+program_source:149:155: error: 'buffer' attribute parameter is out of bounds: must be between 0 and 30
+kernel void main_sort_triangles(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], device Dbg& dbg [[buffer(50)]], device AU32s& counters [[buffer(20)]], device U32s& indices [[buffer(10)]], device F32s& positions [[buffer(11)]], device AI32s& LUT [[buffer(21)]]) {
+ ^
+
diff --git a/test/bug/tint/1113.wgsl.expected.spvasm b/test/bug/tint/1113.wgsl.expected.spvasm
new file mode 100644
index 0000000..fc5ae50
--- /dev/null
+++ b/test/bug/tint/1113.wgsl.expected.spvasm
@@ -0,0 +1,645 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 403
+; Schema: 0
+ OpCapability Shader
+ %65 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main_count "main_count" %GlobalInvocationID_1
+ OpEntryPoint GLCompute %main_create_lut "main_create_lut" %GlobalInvocationID_2
+ OpEntryPoint GLCompute %main_sort_triangles "main_sort_triangles" %GlobalInvocationID_3
+ OpExecutionMode %main_count LocalSize 128 1 1
+ OpExecutionMode %main_create_lut LocalSize 128 1 1
+ OpExecutionMode %main_sort_triangles LocalSize 128 1 1
+ OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
+ OpName %GlobalInvocationID_2 "GlobalInvocationID_2"
+ OpName %GlobalInvocationID_3 "GlobalInvocationID_3"
+ OpName %Uniforms "Uniforms"
+ OpMemberName %Uniforms 0 "numTriangles"
+ OpMemberName %Uniforms 1 "gridSize"
+ OpMemberName %Uniforms 2 "pad1"
+ OpMemberName %Uniforms 3 "pad2"
+ OpMemberName %Uniforms 4 "bbMin"
+ OpMemberName %Uniforms 5 "bbMax"
+ OpName %uniforms "uniforms"
+ OpName %U32s "U32s"
+ OpMemberName %U32s 0 "values"
+ OpName %indices "indices"
+ OpName %F32s "F32s"
+ OpMemberName %F32s 0 "values"
+ OpName %positions "positions"
+ OpName %AU32s "AU32s"
+ OpMemberName %AU32s 0 "values"
+ OpName %counters "counters"
+ OpName %AI32s "AI32s"
+ OpMemberName %AI32s 0 "values"
+ OpName %LUT "LUT"
+ OpName %Dbg "Dbg"
+ OpMemberName %Dbg 0 "offsetCounter"
+ OpMemberName %Dbg 1 "pad0"
+ OpMemberName %Dbg 2 "pad1"
+ OpMemberName %Dbg 3 "pad2"
+ OpMemberName %Dbg 4 "value0"
+ OpMemberName %Dbg 5 "value1"
+ OpMemberName %Dbg 6 "value2"
+ OpMemberName %Dbg 7 "value3"
+ OpMemberName %Dbg 8 "value_f32_0"
+ OpMemberName %Dbg 9 "value_f32_1"
+ OpMemberName %Dbg 10 "value_f32_2"
+ OpMemberName %Dbg 11 "value_f32_3"
+ OpName %dbg "dbg"
+ OpName %toVoxelPos "toVoxelPos"
+ OpName %position "position"
+ OpName %bbMin "bbMin"
+ OpName %bbMax "bbMax"
+ OpName %bbSize "bbSize"
+ OpName %cubeSize "cubeSize"
+ OpName %gridSize "gridSize"
+ OpName %gx "gx"
+ OpName %gy "gy"
+ OpName %gz "gz"
+ OpName %toIndex1D "toIndex1D"
+ OpName %gridSize_0 "gridSize"
+ OpName %voxelPos "voxelPos"
+ OpName %icoord "icoord"
+ OpName %toIndex3D "toIndex3D"
+ OpName %gridSize_1 "gridSize"
+ OpName %index "index"
+ OpName %z "z"
+ OpName %y "y"
+ OpName %x "x"
+ OpName %loadPosition "loadPosition"
+ OpName %vertexIndex "vertexIndex"
+ OpName %position_0 "position"
+ OpName %doIgnore "doIgnore"
+ OpName %g42 "g42"
+ OpName %kj6 "kj6"
+ OpName %b53 "b53"
+ OpName %rwg "rwg"
+ OpName %rb5 "rb5"
+ OpName %g55 "g55"
+ OpName %main_count_inner "main_count_inner"
+ OpName %GlobalInvocationID "GlobalInvocationID"
+ OpName %triangleIndex "triangleIndex"
+ OpName %i0 "i0"
+ OpName %i1 "i1"
+ OpName %i2 "i2"
+ OpName %p0 "p0"
+ OpName %p1 "p1"
+ OpName %p2 "p2"
+ OpName %center "center"
+ OpName %voxelPos_0 "voxelPos"
+ OpName %voxelIndex "voxelIndex"
+ OpName %acefg "acefg"
+ OpName %main_count "main_count"
+ OpName %main_create_lut_inner "main_create_lut_inner"
+ OpName %GlobalInvocationID_0 "GlobalInvocationID"
+ OpName %voxelIndex_0 "voxelIndex"
+ OpName %maxVoxels "maxVoxels"
+ OpName %numTriangles "numTriangles"
+ OpName %offset "offset"
+ OpName %main_create_lut "main_create_lut"
+ OpName %main_sort_triangles_inner "main_sort_triangles_inner"
+ OpName %GlobalInvocationID_4 "GlobalInvocationID"
+ OpName %triangleIndex_0 "triangleIndex"
+ OpName %i0_0 "i0"
+ OpName %i1_0 "i1"
+ OpName %i2_0 "i2"
+ OpName %p0_0 "p0"
+ OpName %p1_0 "p1"
+ OpName %p2_0 "p2"
+ OpName %center_0 "center"
+ OpName %voxelPos_1 "voxelPos"
+ OpName %voxelIndex_1 "voxelIndex"
+ OpName %triangleOffset "triangleOffset"
+ OpName %main_sort_triangles "main_sort_triangles"
+ OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
+ OpDecorate %GlobalInvocationID_2 BuiltIn GlobalInvocationId
+ OpDecorate %GlobalInvocationID_3 BuiltIn GlobalInvocationId
+ OpDecorate %Uniforms Block
+ OpMemberDecorate %Uniforms 0 Offset 0
+ OpMemberDecorate %Uniforms 1 Offset 4
+ OpMemberDecorate %Uniforms 2 Offset 8
+ OpMemberDecorate %Uniforms 3 Offset 12
+ OpMemberDecorate %Uniforms 4 Offset 16
+ OpMemberDecorate %Uniforms 5 Offset 32
+ OpDecorate %uniforms NonWritable
+ OpDecorate %uniforms Binding 0
+ OpDecorate %uniforms DescriptorSet 0
+ OpDecorate %U32s Block
+ OpMemberDecorate %U32s 0 Offset 0
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpDecorate %indices Binding 10
+ OpDecorate %indices DescriptorSet 0
+ OpDecorate %F32s Block
+ OpMemberDecorate %F32s 0 Offset 0
+ OpDecorate %_runtimearr_float ArrayStride 4
+ OpDecorate %positions Binding 11
+ OpDecorate %positions DescriptorSet 0
+ OpDecorate %AU32s Block
+ OpMemberDecorate %AU32s 0 Offset 0
+ OpDecorate %_runtimearr_uint_0 ArrayStride 4
+ OpDecorate %counters Binding 20
+ OpDecorate %counters DescriptorSet 0
+ OpDecorate %AI32s Block
+ OpMemberDecorate %AI32s 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %LUT Binding 21
+ OpDecorate %LUT DescriptorSet 0
+ OpDecorate %Dbg Block
+ OpMemberDecorate %Dbg 0 Offset 0
+ OpMemberDecorate %Dbg 1 Offset 4
+ OpMemberDecorate %Dbg 2 Offset 8
+ OpMemberDecorate %Dbg 3 Offset 12
+ OpMemberDecorate %Dbg 4 Offset 16
+ OpMemberDecorate %Dbg 5 Offset 20
+ OpMemberDecorate %Dbg 6 Offset 24
+ OpMemberDecorate %Dbg 7 Offset 28
+ OpMemberDecorate %Dbg 8 Offset 32
+ OpMemberDecorate %Dbg 9 Offset 36
+ OpMemberDecorate %Dbg 10 Offset 40
+ OpMemberDecorate %Dbg 11 Offset 44
+ OpDecorate %dbg Binding 50
+ OpDecorate %dbg DescriptorSet 0
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
+%GlobalInvocationID_2 = OpVariable %_ptr_Input_v3uint Input
+%GlobalInvocationID_3 = OpVariable %_ptr_Input_v3uint Input
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %Uniforms = OpTypeStruct %uint %uint %uint %uint %v3float %v3float
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+ %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %U32s = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer_U32s = OpTypePointer StorageBuffer %U32s
+ %indices = OpVariable %_ptr_StorageBuffer_U32s StorageBuffer
+%_runtimearr_float = OpTypeRuntimeArray %float
+ %F32s = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_F32s = OpTypePointer StorageBuffer %F32s
+ %positions = OpVariable %_ptr_StorageBuffer_F32s StorageBuffer
+%_runtimearr_uint_0 = OpTypeRuntimeArray %uint
+ %AU32s = OpTypeStruct %_runtimearr_uint_0
+%_ptr_StorageBuffer_AU32s = OpTypePointer StorageBuffer %AU32s
+ %counters = OpVariable %_ptr_StorageBuffer_AU32s StorageBuffer
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %AI32s = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_AI32s = OpTypePointer StorageBuffer %AI32s
+ %LUT = OpVariable %_ptr_StorageBuffer_AI32s StorageBuffer
+ %Dbg = OpTypeStruct %uint %uint %uint %uint %uint %uint %uint %uint %float %float %float %float
+%_ptr_StorageBuffer_Dbg = OpTypePointer StorageBuffer %Dbg
+ %dbg = OpVariable %_ptr_StorageBuffer_Dbg StorageBuffer
+ %32 = OpTypeFunction %v3float %v3float
+ %uint_4 = OpConstant %uint 4
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %50 = OpConstantNull %v3float
+ %uint_5 = OpConstant %uint 5
+%_ptr_Function_float = OpTypePointer Function %float
+ %75 = OpConstantNull %float
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %112 = OpTypeFunction %uint %uint %v3float
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+ %120 = OpConstantNull %v3uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %133 = OpTypeFunction %v3uint %uint %uint
+ %141 = OpConstantNull %uint
+ %154 = OpTypeFunction %v3float %uint
+ %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %void = OpTypeVoid
+ %175 = OpTypeFunction %void
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_uint_0 = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %204 = OpConstantNull %int
+ %205 = OpTypeFunction %void %v3uint
+ %bool = OpTypeBool
+ %float_3 = OpConstant %float 3
+ %uint_8 = OpConstant %uint 8
+ %uint_9 = OpConstant %uint 9
+ %uint_10 = OpConstant %uint 10
+ %int_n1 = OpConstant %int -1
+ %int_1 = OpConstant %int 1
+ %toVoxelPos = OpFunction %v3float None %32
+ %position = OpFunctionParameter %v3float
+ %35 = OpLabel
+ %bbMin = OpVariable %_ptr_Function_v3float Function %50
+ %bbMax = OpVariable %_ptr_Function_v3float Function %50
+ %bbSize = OpVariable %_ptr_Function_v3float Function %50
+ %cubeSize = OpVariable %_ptr_Function_float Function %75
+ %gridSize = OpVariable %_ptr_Function_float Function %75
+ %gx = OpVariable %_ptr_Function_float Function %75
+ %gy = OpVariable %_ptr_Function_float Function %75
+ %gz = OpVariable %_ptr_Function_float Function %75
+ %39 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_0
+ %40 = OpLoad %float %39
+ %42 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_1
+ %43 = OpLoad %float %42
+ %45 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_2
+ %46 = OpLoad %float %45
+ %47 = OpCompositeConstruct %v3float %40 %43 %46
+ OpStore %bbMin %47
+ %52 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_5 %uint_0
+ %53 = OpLoad %float %52
+ %54 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_5 %uint_1
+ %55 = OpLoad %float %54
+ %56 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_5 %uint_2
+ %57 = OpLoad %float %56
+ %58 = OpCompositeConstruct %v3float %53 %55 %57
+ OpStore %bbMax %58
+ %60 = OpLoad %v3float %bbMax
+ %61 = OpLoad %v3float %bbMin
+ %62 = OpFSub %v3float %60 %61
+ OpStore %bbSize %62
+ %68 = OpAccessChain %_ptr_Function_float %bbSize %uint_0
+ %69 = OpLoad %float %68
+ %70 = OpAccessChain %_ptr_Function_float %bbSize %uint_1
+ %71 = OpLoad %float %70
+ %66 = OpExtInst %float %65 NMax %69 %71
+ %72 = OpAccessChain %_ptr_Function_float %bbSize %uint_2
+ %73 = OpLoad %float %72
+ %64 = OpExtInst %float %65 NMax %66 %73
+ OpStore %cubeSize %64
+ %78 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %79 = OpLoad %uint %78
+ %76 = OpConvertUToF %float %79
+ OpStore %gridSize %76
+ %81 = OpLoad %float %gridSize
+ %82 = OpCompositeExtract %float %position 0
+ %83 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_0
+ %84 = OpLoad %float %83
+ %85 = OpFSub %float %82 %84
+ %86 = OpFMul %float %81 %85
+ %87 = OpLoad %float %cubeSize
+ %88 = OpFDiv %float %86 %87
+ OpStore %gx %88
+ %90 = OpLoad %float %gridSize
+ %91 = OpCompositeExtract %float %position 1
+ %92 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_1
+ %93 = OpLoad %float %92
+ %94 = OpFSub %float %91 %93
+ %95 = OpFMul %float %90 %94
+ %96 = OpLoad %float %cubeSize
+ %97 = OpFDiv %float %95 %96
+ OpStore %gy %97
+ %99 = OpLoad %float %gridSize
+ %100 = OpCompositeExtract %float %position 2
+ %101 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_4 %uint_2
+ %102 = OpLoad %float %101
+ %103 = OpFSub %float %100 %102
+ %104 = OpFMul %float %99 %103
+ %105 = OpLoad %float %cubeSize
+ %106 = OpFDiv %float %104 %105
+ OpStore %gz %106
+ %108 = OpLoad %float %gx
+ %109 = OpLoad %float %gy
+ %110 = OpLoad %float %gz
+ %111 = OpCompositeConstruct %v3float %108 %109 %110
+ OpReturnValue %111
+ OpFunctionEnd
+ %toIndex1D = OpFunction %uint None %112
+ %gridSize_0 = OpFunctionParameter %uint
+ %voxelPos = OpFunctionParameter %v3float
+ %116 = OpLabel
+ %icoord = OpVariable %_ptr_Function_v3uint Function %120
+ %117 = OpConvertFToU %v3uint %voxelPos
+ OpStore %icoord %117
+ %122 = OpAccessChain %_ptr_Function_uint %icoord %uint_0
+ %123 = OpLoad %uint %122
+ %124 = OpAccessChain %_ptr_Function_uint %icoord %uint_1
+ %125 = OpLoad %uint %124
+ %126 = OpIMul %uint %gridSize_0 %125
+ %127 = OpIAdd %uint %123 %126
+ %128 = OpIMul %uint %gridSize_0 %gridSize_0
+ %129 = OpAccessChain %_ptr_Function_uint %icoord %uint_2
+ %130 = OpLoad %uint %129
+ %131 = OpIMul %uint %128 %130
+ %132 = OpIAdd %uint %127 %131
+ OpReturnValue %132
+ OpFunctionEnd
+ %toIndex3D = OpFunction %v3uint None %133
+ %gridSize_1 = OpFunctionParameter %uint
+ %index = OpFunctionParameter %uint
+ %137 = OpLabel
+ %z = OpVariable %_ptr_Function_uint Function %141
+ %y = OpVariable %_ptr_Function_uint Function %141
+ %x = OpVariable %_ptr_Function_uint Function %141
+ %138 = OpIMul %uint %gridSize_1 %gridSize_1
+ %139 = OpUDiv %uint %index %138
+ OpStore %z %139
+ %142 = OpIMul %uint %gridSize_1 %gridSize_1
+ %143 = OpLoad %uint %z
+ %144 = OpIMul %uint %142 %143
+ %145 = OpISub %uint %index %144
+ %146 = OpUDiv %uint %145 %gridSize_1
+ OpStore %y %146
+ %148 = OpUMod %uint %index %gridSize_1
+ OpStore %x %148
+ %150 = OpLoad %uint %x
+ %151 = OpLoad %uint %y
+ %152 = OpLoad %uint %z
+ %153 = OpCompositeConstruct %v3uint %150 %151 %152
+ OpReturnValue %153
+ OpFunctionEnd
+%loadPosition = OpFunction %v3float None %154
+%vertexIndex = OpFunctionParameter %uint
+ %157 = OpLabel
+ %position_0 = OpVariable %_ptr_Function_v3float Function %50
+ %159 = OpIMul %uint %uint_3 %vertexIndex
+ %160 = OpIAdd %uint %159 %uint_0
+ %162 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %160
+ %163 = OpLoad %float %162
+ %164 = OpIMul %uint %uint_3 %vertexIndex
+ %165 = OpIAdd %uint %164 %uint_1
+ %166 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %165
+ %167 = OpLoad %float %166
+ %168 = OpIMul %uint %uint_3 %vertexIndex
+ %169 = OpIAdd %uint %168 %uint_2
+ %170 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %169
+ %171 = OpLoad %float %170
+ %172 = OpCompositeConstruct %v3float %163 %167 %171
+ OpStore %position_0 %172
+ %174 = OpLoad %v3float %position_0
+ OpReturnValue %174
+ OpFunctionEnd
+ %doIgnore = OpFunction %void None %175
+ %178 = OpLabel
+ %g42 = OpVariable %_ptr_Function_uint Function %141
+ %kj6 = OpVariable %_ptr_Function_uint Function %141
+ %b53 = OpVariable %_ptr_Function_uint Function %141
+ %rwg = OpVariable %_ptr_Function_uint Function %141
+ %rb5 = OpVariable %_ptr_Function_float Function %75
+ %g55 = OpVariable %_ptr_Function_int Function %204
+ %179 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+ %180 = OpLoad %uint %179
+ OpStore %g42 %180
+ %183 = OpAccessChain %_ptr_StorageBuffer_uint %dbg %uint_5
+ %184 = OpLoad %uint %183
+ OpStore %kj6 %184
+ %190 = OpAccessChain %_ptr_StorageBuffer_uint_0 %counters %uint_0 %int_0
+ %186 = OpAtomicLoad %uint %190 %uint_1 %uint_0
+ OpStore %b53 %186
+ %192 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %int_0
+ %193 = OpLoad %uint %192
+ OpStore %rwg %193
+ %195 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %int_0
+ %196 = OpLoad %float %195
+ OpStore %rb5 %196
+ %201 = OpAccessChain %_ptr_StorageBuffer_int %LUT %uint_0 %int_0
+ %198 = OpAtomicLoad %int %201 %uint_1 %uint_0
+ OpStore %g55 %198
+ OpReturn
+ OpFunctionEnd
+%main_count_inner = OpFunction %void None %205
+%GlobalInvocationID = OpFunctionParameter %v3uint
+ %208 = OpLabel
+%triangleIndex = OpVariable %_ptr_Function_uint Function %141
+ %i0 = OpVariable %_ptr_Function_uint Function %141
+ %i1 = OpVariable %_ptr_Function_uint Function %141
+ %i2 = OpVariable %_ptr_Function_uint Function %141
+ %p0 = OpVariable %_ptr_Function_v3float Function %50
+ %p1 = OpVariable %_ptr_Function_v3float Function %50
+ %p2 = OpVariable %_ptr_Function_v3float Function %50
+ %253 = OpVariable %_ptr_Function_v3float Function %50
+ %center = OpVariable %_ptr_Function_v3float Function %50
+ %voxelPos_0 = OpVariable %_ptr_Function_v3float Function %50
+ %voxelIndex = OpVariable %_ptr_Function_uint Function %141
+ %acefg = OpVariable %_ptr_Function_uint Function %141
+ %209 = OpCompositeExtract %uint %GlobalInvocationID 0
+ OpStore %triangleIndex %209
+ %211 = OpLoad %uint %triangleIndex
+ %212 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+ %213 = OpLoad %uint %212
+ %214 = OpUGreaterThanEqual %bool %211 %213
+ OpSelectionMerge %216 None
+ OpBranchConditional %214 %217 %216
+ %217 = OpLabel
+ OpReturn
+ %216 = OpLabel
+ %218 = OpFunctionCall %void %doIgnore
+ %219 = OpLoad %uint %triangleIndex
+ %220 = OpIMul %uint %uint_3 %219
+ %221 = OpIAdd %uint %220 %uint_0
+ %222 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %221
+ %223 = OpLoad %uint %222
+ OpStore %i0 %223
+ %225 = OpLoad %uint %triangleIndex
+ %226 = OpIMul %uint %uint_3 %225
+ %227 = OpIAdd %uint %226 %uint_1
+ %228 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %227
+ %229 = OpLoad %uint %228
+ OpStore %i1 %229
+ %231 = OpLoad %uint %triangleIndex
+ %232 = OpIMul %uint %uint_3 %231
+ %233 = OpIAdd %uint %232 %uint_2
+ %234 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %233
+ %235 = OpLoad %uint %234
+ OpStore %i2 %235
+ %238 = OpLoad %uint %i0
+ %237 = OpFunctionCall %v3float %loadPosition %238
+ OpStore %p0 %237
+ %241 = OpLoad %uint %i1
+ %240 = OpFunctionCall %v3float %loadPosition %241
+ OpStore %p1 %240
+ %244 = OpLoad %uint %i2
+ %243 = OpFunctionCall %v3float %loadPosition %244
+ OpStore %p2 %243
+ %246 = OpLoad %v3float %p0
+ %247 = OpLoad %v3float %p1
+ %248 = OpFAdd %v3float %246 %247
+ %249 = OpLoad %v3float %p2
+ %250 = OpFAdd %v3float %248 %249
+ %254 = OpCompositeConstruct %v3float %float_3 %float_3 %float_3
+ %252 = OpFDiv %v3float %250 %254
+ OpStore %center %252
+ %257 = OpLoad %v3float %center
+ %256 = OpFunctionCall %v3float %toVoxelPos %257
+ OpStore %voxelPos_0 %256
+ %260 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %261 = OpLoad %uint %260
+ %262 = OpLoad %v3float %voxelPos_0
+ %259 = OpFunctionCall %uint %toIndex1D %261 %262
+ OpStore %voxelIndex %259
+ %266 = OpLoad %uint %voxelIndex
+ %267 = OpAccessChain %_ptr_StorageBuffer_uint_0 %counters %uint_0 %266
+ %264 = OpAtomicIAdd %uint %267 %uint_1 %uint_0 %uint_1
+ OpStore %acefg %264
+ %269 = OpLoad %uint %triangleIndex
+ %270 = OpIEqual %bool %269 %uint_0
+ OpSelectionMerge %271 None
+ OpBranchConditional %270 %272 %271
+ %272 = OpLabel
+ %273 = OpAccessChain %_ptr_StorageBuffer_uint %dbg %uint_4
+ %274 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %275 = OpLoad %uint %274
+ OpStore %273 %275
+ %277 = OpAccessChain %_ptr_StorageBuffer_float %dbg %uint_8
+ %278 = OpAccessChain %_ptr_Function_float %center %uint_0
+ %279 = OpLoad %float %278
+ OpStore %277 %279
+ %281 = OpAccessChain %_ptr_StorageBuffer_float %dbg %uint_9
+ %282 = OpAccessChain %_ptr_Function_float %center %uint_1
+ %283 = OpLoad %float %282
+ OpStore %281 %283
+ %285 = OpAccessChain %_ptr_StorageBuffer_float %dbg %uint_10
+ %286 = OpAccessChain %_ptr_Function_float %center %uint_2
+ %287 = OpLoad %float %286
+ OpStore %285 %287
+ OpBranch %271
+ %271 = OpLabel
+ OpReturn
+ OpFunctionEnd
+ %main_count = OpFunction %void None %175
+ %289 = OpLabel
+ %291 = OpLoad %v3uint %GlobalInvocationID_1
+ %290 = OpFunctionCall %void %main_count_inner %291
+ OpReturn
+ OpFunctionEnd
+%main_create_lut_inner = OpFunction %void None %205
+%GlobalInvocationID_0 = OpFunctionParameter %v3uint
+ %294 = OpLabel
+%voxelIndex_0 = OpVariable %_ptr_Function_uint Function %141
+ %maxVoxels = OpVariable %_ptr_Function_uint Function %141
+%numTriangles = OpVariable %_ptr_Function_uint Function %141
+ %offset = OpVariable %_ptr_Function_int Function %204
+ %295 = OpCompositeExtract %uint %GlobalInvocationID_0 0
+ OpStore %voxelIndex_0 %295
+ %297 = OpFunctionCall %void %doIgnore
+ %298 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %299 = OpLoad %uint %298
+ %300 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %301 = OpLoad %uint %300
+ %302 = OpIMul %uint %299 %301
+ %303 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %304 = OpLoad %uint %303
+ %305 = OpIMul %uint %302 %304
+ OpStore %maxVoxels %305
+ %307 = OpLoad %uint %voxelIndex_0
+ %308 = OpLoad %uint %maxVoxels
+ %309 = OpUGreaterThanEqual %bool %307 %308
+ OpSelectionMerge %310 None
+ OpBranchConditional %309 %311 %310
+ %311 = OpLabel
+ OpReturn
+ %310 = OpLabel
+ %314 = OpLoad %uint %voxelIndex_0
+ %315 = OpAccessChain %_ptr_StorageBuffer_uint_0 %counters %uint_0 %314
+ %312 = OpAtomicLoad %uint %315 %uint_1 %uint_0
+ OpStore %numTriangles %312
+ OpStore %offset %int_n1
+ %319 = OpLoad %uint %numTriangles
+ %320 = OpUGreaterThan %bool %319 %uint_0
+ OpSelectionMerge %321 None
+ OpBranchConditional %320 %322 %321
+ %322 = OpLabel
+ %326 = OpAccessChain %_ptr_StorageBuffer_uint_0 %dbg %uint_0
+ %327 = OpLoad %uint %numTriangles
+ %324 = OpAtomicIAdd %uint %326 %uint_1 %uint_0 %327
+ %323 = OpBitcast %int %324
+ OpStore %offset %323
+ OpBranch %321
+ %321 = OpLabel
+ %330 = OpLoad %uint %voxelIndex_0
+ %331 = OpAccessChain %_ptr_StorageBuffer_int %LUT %uint_0 %330
+ %332 = OpLoad %int %offset
+ OpAtomicStore %331 %uint_1 %uint_0 %332
+ OpReturn
+ OpFunctionEnd
+%main_create_lut = OpFunction %void None %175
+ %334 = OpLabel
+ %336 = OpLoad %v3uint %GlobalInvocationID_2
+ %335 = OpFunctionCall %void %main_create_lut_inner %336
+ OpReturn
+ OpFunctionEnd
+%main_sort_triangles_inner = OpFunction %void None %205
+%GlobalInvocationID_4 = OpFunctionParameter %v3uint
+ %339 = OpLabel
+%triangleIndex_0 = OpVariable %_ptr_Function_uint Function %141
+ %i0_0 = OpVariable %_ptr_Function_uint Function %141
+ %i1_0 = OpVariable %_ptr_Function_uint Function %141
+ %i2_0 = OpVariable %_ptr_Function_uint Function %141
+ %p0_0 = OpVariable %_ptr_Function_v3float Function %50
+ %p1_0 = OpVariable %_ptr_Function_v3float Function %50
+ %p2_0 = OpVariable %_ptr_Function_v3float Function %50
+ %382 = OpVariable %_ptr_Function_v3float Function %50
+ %center_0 = OpVariable %_ptr_Function_v3float Function %50
+ %voxelPos_1 = OpVariable %_ptr_Function_v3float Function %50
+%voxelIndex_1 = OpVariable %_ptr_Function_uint Function %141
+%triangleOffset = OpVariable %_ptr_Function_int Function %204
+ %340 = OpCompositeExtract %uint %GlobalInvocationID_4 0
+ OpStore %triangleIndex_0 %340
+ %342 = OpFunctionCall %void %doIgnore
+ %343 = OpLoad %uint %triangleIndex_0
+ %344 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+ %345 = OpLoad %uint %344
+ %346 = OpUGreaterThanEqual %bool %343 %345
+ OpSelectionMerge %347 None
+ OpBranchConditional %346 %348 %347
+ %348 = OpLabel
+ OpReturn
+ %347 = OpLabel
+ %349 = OpLoad %uint %triangleIndex_0
+ %350 = OpIMul %uint %uint_3 %349
+ %351 = OpIAdd %uint %350 %uint_0
+ %352 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %351
+ %353 = OpLoad %uint %352
+ OpStore %i0_0 %353
+ %355 = OpLoad %uint %triangleIndex_0
+ %356 = OpIMul %uint %uint_3 %355
+ %357 = OpIAdd %uint %356 %uint_1
+ %358 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %357
+ %359 = OpLoad %uint %358
+ OpStore %i1_0 %359
+ %361 = OpLoad %uint %triangleIndex_0
+ %362 = OpIMul %uint %uint_3 %361
+ %363 = OpIAdd %uint %362 %uint_2
+ %364 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %363
+ %365 = OpLoad %uint %364
+ OpStore %i2_0 %365
+ %368 = OpLoad %uint %i0_0
+ %367 = OpFunctionCall %v3float %loadPosition %368
+ OpStore %p0_0 %367
+ %371 = OpLoad %uint %i1_0
+ %370 = OpFunctionCall %v3float %loadPosition %371
+ OpStore %p1_0 %370
+ %374 = OpLoad %uint %i2_0
+ %373 = OpFunctionCall %v3float %loadPosition %374
+ OpStore %p2_0 %373
+ %376 = OpLoad %v3float %p0_0
+ %377 = OpLoad %v3float %p1_0
+ %378 = OpFAdd %v3float %376 %377
+ %379 = OpLoad %v3float %p2_0
+ %380 = OpFAdd %v3float %378 %379
+ %383 = OpCompositeConstruct %v3float %float_3 %float_3 %float_3
+ %381 = OpFDiv %v3float %380 %383
+ OpStore %center_0 %381
+ %386 = OpLoad %v3float %center_0
+ %385 = OpFunctionCall %v3float %toVoxelPos %386
+ OpStore %voxelPos_1 %385
+ %389 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+ %390 = OpLoad %uint %389
+ %391 = OpLoad %v3float %voxelPos_1
+ %388 = OpFunctionCall %uint %toIndex1D %390 %391
+ OpStore %voxelIndex_1 %388
+ %395 = OpLoad %uint %voxelIndex_1
+ %396 = OpAccessChain %_ptr_StorageBuffer_int %LUT %uint_0 %395
+ %393 = OpAtomicIAdd %int %396 %uint_1 %uint_0 %int_1
+ OpStore %triangleOffset %393
+ OpReturn
+ OpFunctionEnd
+%main_sort_triangles = OpFunction %void None %175
+ %400 = OpLabel
+ %402 = OpLoad %v3uint %GlobalInvocationID_3
+ %401 = OpFunctionCall %void %main_sort_triangles_inner %402
+ OpReturn
+ OpFunctionEnd
diff --git a/test/bug/tint/1113.wgsl.expected.wgsl b/test/bug/tint/1113.wgsl.expected.wgsl
new file mode 100644
index 0000000..cd39fdf
--- /dev/null
+++ b/test/bug/tint/1113.wgsl.expected.wgsl
@@ -0,0 +1,160 @@
+[[block]]
+struct Uniforms {
+ numTriangles : u32;
+ gridSize : u32;
+ pad1 : u32;
+ pad2 : u32;
+ bbMin : vec3<f32>;
+ bbMax : vec3<f32>;
+};
+
+[[block]]
+struct Dbg {
+ offsetCounter : atomic<u32>;
+ pad0 : u32;
+ pad1 : u32;
+ pad2 : u32;
+ value0 : u32;
+ value1 : u32;
+ value2 : u32;
+ value3 : u32;
+ value_f32_0 : f32;
+ value_f32_1 : f32;
+ value_f32_2 : f32;
+ value_f32_3 : f32;
+};
+
+[[block]]
+struct F32s {
+ values : [[stride(4)]] array<f32>;
+};
+
+[[block]]
+struct U32s {
+ values : [[stride(4)]] array<u32>;
+};
+
+[[block]]
+struct I32s {
+ values : [[stride(4)]] array<i32>;
+};
+
+[[block]]
+struct AU32s {
+ values : [[stride(4)]] array<atomic<u32>>;
+};
+
+[[block]]
+struct AI32s {
+ values : [[stride(4)]] array<atomic<i32>>;
+};
+
+[[binding(0), group(0)]] var<uniform> uniforms : Uniforms;
+
+[[binding(10), group(0)]] var<storage, read_write> indices : U32s;
+
+[[binding(11), group(0)]] var<storage, read_write> positions : F32s;
+
+[[binding(20), group(0)]] var<storage, read_write> counters : AU32s;
+
+[[binding(21), group(0)]] var<storage, read_write> LUT : AI32s;
+
+[[binding(50), group(0)]] var<storage, read_write> dbg : Dbg;
+
+fn toVoxelPos(position : vec3<f32>) -> vec3<f32> {
+ var bbMin = vec3<f32>(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z);
+ var bbMax = vec3<f32>(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z);
+ var bbSize = (bbMax - bbMin);
+ var cubeSize = max(max(bbSize.x, bbSize.y), bbSize.z);
+ var gridSize = f32(uniforms.gridSize);
+ var gx = ((gridSize * (position.x - uniforms.bbMin.x)) / cubeSize);
+ var gy = ((gridSize * (position.y - uniforms.bbMin.y)) / cubeSize);
+ var gz = ((gridSize * (position.z - uniforms.bbMin.z)) / cubeSize);
+ return vec3<f32>(gx, gy, gz);
+}
+
+fn toIndex1D(gridSize : u32, voxelPos : vec3<f32>) -> u32 {
+ var icoord = vec3<u32>(voxelPos);
+ return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
+}
+
+fn toIndex3D(gridSize : u32, index : u32) -> vec3<u32> {
+ var z = (index / (gridSize * gridSize));
+ var y = ((index - ((gridSize * gridSize) * z)) / gridSize);
+ var x = (index % gridSize);
+ return vec3<u32>(x, y, z);
+}
+
+fn loadPosition(vertexIndex : u32) -> vec3<f32> {
+ var position = vec3<f32>(positions.values[((3u * vertexIndex) + 0u)], positions.values[((3u * vertexIndex) + 1u)], positions.values[((3u * vertexIndex) + 2u)]);
+ return position;
+}
+
+fn doIgnore() {
+ var g42 = uniforms.numTriangles;
+ var kj6 = dbg.value1;
+ var b53 = atomicLoad(&(counters.values[0]));
+ var rwg = indices.values[0];
+ var rb5 = positions.values[0];
+ var g55 = atomicLoad(&(LUT.values[0]));
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_count([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ var triangleIndex = GlobalInvocationID.x;
+ if ((triangleIndex >= uniforms.numTriangles)) {
+ return;
+ }
+ doIgnore();
+ var i0 = indices.values[((3u * triangleIndex) + 0u)];
+ var i1 = indices.values[((3u * triangleIndex) + 1u)];
+ var i2 = indices.values[((3u * triangleIndex) + 2u)];
+ var p0 = loadPosition(i0);
+ var p1 = loadPosition(i1);
+ var p2 = loadPosition(i2);
+ var center = (((p0 + p1) + p2) / 3.0);
+ var voxelPos = toVoxelPos(center);
+ var voxelIndex = toIndex1D(uniforms.gridSize, voxelPos);
+ var acefg = atomicAdd(&(counters.values[voxelIndex]), 1u);
+ if ((triangleIndex == 0u)) {
+ dbg.value0 = uniforms.gridSize;
+ dbg.value_f32_0 = center.x;
+ dbg.value_f32_1 = center.y;
+ dbg.value_f32_2 = center.z;
+ }
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_create_lut([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ var voxelIndex = GlobalInvocationID.x;
+ doIgnore();
+ var maxVoxels = ((uniforms.gridSize * uniforms.gridSize) * uniforms.gridSize);
+ if ((voxelIndex >= maxVoxels)) {
+ return;
+ }
+ var numTriangles = atomicLoad(&(counters.values[voxelIndex]));
+ var offset = -1;
+ if ((numTriangles > 0u)) {
+ offset = i32(atomicAdd(&(dbg.offsetCounter), numTriangles));
+ }
+ atomicStore(&(LUT.values[voxelIndex]), offset);
+}
+
+[[stage(compute), workgroup_size(128)]]
+fn main_sort_triangles([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+ var triangleIndex = GlobalInvocationID.x;
+ doIgnore();
+ if ((triangleIndex >= uniforms.numTriangles)) {
+ return;
+ }
+ var i0 = indices.values[((3u * triangleIndex) + 0u)];
+ var i1 = indices.values[((3u * triangleIndex) + 1u)];
+ var i2 = indices.values[((3u * triangleIndex) + 2u)];
+ var p0 = loadPosition(i0);
+ var p1 = loadPosition(i1);
+ var p2 = loadPosition(i2);
+ var center = (((p0 + p1) + p2) / 3.0);
+ var voxelPos = toVoxelPos(center);
+ var voxelIndex = toIndex1D(uniforms.gridSize, voxelPos);
+ var triangleOffset = atomicAdd(&(LUT.values[voxelIndex]), 1);
+}
