reader/wgsl: Generate ForLoopStatements
Instead of LoopStatements.
Update the writers to handle these.
Fixed: tint:952
Change-Id: Ibef66e133224810efc28c224d910b5e21f71f8d6
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57203
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl
index 2c9ea16..e7a735c 100644
--- a/test/bug/tint/534.wgsl.expected.hlsl
+++ b/test/bug/tint/534.wgsl.expected.hlsl
@@ -43,7 +43,7 @@
uint i = 0u;
while (true) {
const uint scalar_offset_1 = (12u) / 4;
- if (!(!(!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))))) { break; }
+ if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) { break; }
Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
bool tint_tmp_1 = success;
if (tint_tmp_1) {
diff --git a/test/bug/tint/534.wgsl.expected.msl b/test/bug/tint/534.wgsl.expected.msl
index 429c94b..356f44d 100644
--- a/test/bug/tint/534.wgsl.expected.msl
+++ b/test/bug/tint/534.wgsl.expected.msl
@@ -27,18 +27,9 @@
bool success = true;
uint4 srcColorBits = 0u;
uint4 dstColorBits = uint4(dstColor);
- {
- uint i = 0u;
- while (true) {
- if (!((i < uniforms.channelCount))) {
- break;
- }
- srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
- success = (success && (srcColorBits[i] == dstColorBits[i]));
- {
- i = (i + 1u);
- }
- }
+ for(uint i = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
+ srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+ success = (success && (srcColorBits[i] == dstColorBits[i]));
}
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
if (success) {
diff --git a/test/bug/tint/534.wgsl.expected.wgsl b/test/bug/tint/534.wgsl.expected.wgsl
index 91a330b..1fd409c 100644
--- a/test/bug/tint/534.wgsl.expected.wgsl
+++ b/test/bug/tint/534.wgsl.expected.wgsl
@@ -36,19 +36,9 @@
var success : bool = true;
var srcColorBits : vec4<u32>;
var dstColorBits : vec4<u32> = vec4<u32>(dstColor);
- {
- var i : u32 = 0u;
- loop {
- if (!((i < uniforms.channelCount))) {
- break;
- }
- srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
- success = (success && (srcColorBits[i] == dstColorBits[i]));
-
- continuing {
- i = (i + 1u);
- }
- }
+ for(var i : u32 = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
+ srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+ success = (success && (srcColorBits[i] == dstColorBits[i]));
}
var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x);
if (success) {
diff --git a/test/bug/tint/744.wgsl.expected.hlsl b/test/bug/tint/744.wgsl.expected.hlsl
index 0314864..47f47a7 100644
--- a/test/bug/tint/744.wgsl.expected.hlsl
+++ b/test/bug/tint/744.wgsl.expected.hlsl
@@ -19,8 +19,7 @@
const uint dimOutter = uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4];
uint result = 0u;
{
- uint i = 0u;
- for(; !(!((i < dimInner))); i = (i + 1u)) {
+ for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
const uint a = (i + (resultCell.x * dimInner));
const uint b = (resultCell.y + (i * dimOutter));
result = (result + (firstMatrix.Load((4u * a)) * secondMatrix.Load((4u * b))));
diff --git a/test/bug/tint/744.wgsl.expected.msl b/test/bug/tint/744.wgsl.expected.msl
index 7855d18..16ebd85 100644
--- a/test/bug/tint/744.wgsl.expected.msl
+++ b/test/bug/tint/744.wgsl.expected.msl
@@ -15,19 +15,10 @@
uint const dimInner = uniforms.aShape.y;
uint const dimOutter = uniforms.outShape.y;
uint result = 0u;
- {
- uint i = 0u;
- while (true) {
- if (!((i < dimInner))) {
- break;
- }
- uint const a = (i + (resultCell.x * dimInner));
- uint const b = (resultCell.y + (i * dimOutter));
- result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
- {
- i = (i + 1u);
- }
- }
+ for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
+ uint const a = (i + (resultCell.x * dimInner));
+ uint const b = (resultCell.y + (i * dimOutter));
+ result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
}
uint const index = (resultCell.y + (resultCell.x * dimOutter));
resultMatrix.numbers[index] = result;
diff --git a/test/bug/tint/744.wgsl.expected.wgsl b/test/bug/tint/744.wgsl.expected.wgsl
index d96d510..ad6c01e 100644
--- a/test/bug/tint/744.wgsl.expected.wgsl
+++ b/test/bug/tint/744.wgsl.expected.wgsl
@@ -24,20 +24,10 @@
let dimInner : u32 = uniforms.aShape.y;
let dimOutter : u32 = uniforms.outShape.y;
var result : u32 = 0u;
- {
- var i : u32 = 0u;
- loop {
- if (!((i < dimInner))) {
- break;
- }
- let a : u32 = (i + (resultCell.x * dimInner));
- let b : u32 = (resultCell.y + (i * dimOutter));
- result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
-
- continuing {
- i = (i + 1u);
- }
- }
+ for(var i : u32 = 0u; (i < dimInner); i = (i + 1u)) {
+ let a : u32 = (i + (resultCell.x * dimInner));
+ let b : u32 = (resultCell.y + (i * dimOutter));
+ result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
}
let index : u32 = (resultCell.y + (resultCell.x * dimOutter));
resultMatrix.numbers[index] = result;
diff --git a/test/bug/tint/757.wgsl.expected.hlsl b/test/bug/tint/757.wgsl.expected.hlsl
index 72ed5bd..da3073c 100644
--- a/test/bug/tint/757.wgsl.expected.hlsl
+++ b/test/bug/tint/757.wgsl.expected.hlsl
@@ -16,8 +16,7 @@
flatIndex = (flatIndex * 1u);
float4 texel = myTexture.Load(int4(GlobalInvocationID.xy, 0, 0));
{
- uint i = 0u;
- for(; !(!((i < 1u))); i = (i + 1u)) {
+ for(uint i = 0u; (i < 1u); i = (i + 1u)) {
result.Store((4u * (flatIndex + i)), asuint(texel.r));
}
}
diff --git a/test/bug/tint/757.wgsl.expected.msl b/test/bug/tint/757.wgsl.expected.msl
index 03e2067..8699607 100644
--- a/test/bug/tint/757.wgsl.expected.msl
+++ b/test/bug/tint/757.wgsl.expected.msl
@@ -12,17 +12,8 @@
uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
flatIndex = (flatIndex * 1u);
float4 texel = tint_symbol_2.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
- {
- uint i = 0u;
- while (true) {
- if (!((i < 1u))) {
- break;
- }
- result.values[(flatIndex + i)] = texel.r;
- {
- i = (i + 1u);
- }
- }
+ for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+ result.values[(flatIndex + i)] = texel.r;
}
return;
}
diff --git a/test/bug/tint/757.wgsl.expected.wgsl b/test/bug/tint/757.wgsl.expected.wgsl
index c7c0d8c..9690af5 100644
--- a/test/bug/tint/757.wgsl.expected.wgsl
+++ b/test/bug/tint/757.wgsl.expected.wgsl
@@ -19,17 +19,7 @@
var flatIndex : u32 = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
flatIndex = (flatIndex * 1u);
var texel : vec4<f32> = textureLoad(myTexture, vec2<i32>(GlobalInvocationID.xy), 0, 0);
- {
- var i : u32 = 0u;
- loop {
- if (!((i < 1u))) {
- break;
- }
- result.values[(flatIndex + i)] = texel.r;
-
- continuing {
- i = (i + 1u);
- }
- }
+ for(var i : u32 = 0u; (i < 1u); i = (i + 1u)) {
+ result.values[(flatIndex + i)] = texel.r;
}
}
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 3de9408..5e0ebb4 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -69,14 +69,22 @@
const uint3 global_id = tint_symbol.global_id;
const uint local_invocation_index = tint_symbol.local_invocation_index;
if ((local_invocation_index == 0u)) {
- for(int i = 0; (i < 64); i = (i + 1)) {
- for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
- mm_Asub[i][i_1] = 0.0f;
+ {
+ for(int i = 0; (i < 64); i = (i + 1)) {
+ {
+ for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+ mm_Asub[i][i_1] = 0.0f;
+ }
+ }
}
}
- for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
- for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
- mm_Bsub[i_2][i_3] = 0.0f;
+ {
+ for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+ {
+ for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
+ mm_Bsub[i_2][i_3] = 0.0f;
+ }
+ }
}
}
}
@@ -91,8 +99,7 @@
float ACached = 0.0f;
float BCached[4] = (float[4])0;
{
- uint index = 0u;
- for(; !(!((index < (RowPerThread * ColPerThread)))); index = (index + 1u)) {
+ for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
acc[index] = 0.0f;
}
}
@@ -101,14 +108,11 @@
const uint RowPerThreadB = (TileInner / 16u);
const uint tileRowB = (local_id.y * RowPerThreadB);
{
- uint t = 0u;
- for(; !(!((t < numTiles))); t = (t + 1u)) {
+ for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThreadA))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRow + innerRow);
const uint inputCol = (tileColA + innerCol);
mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
@@ -117,11 +121,9 @@
}
}
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThreadB))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRowB + innerRow);
const uint inputCol = (tileCol + innerCol);
mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
@@ -131,21 +133,17 @@
}
GroupMemoryBarrierWithGroupSync();
{
- uint k = 0u;
- for(; !(!((k < TileInner))); k = (k + 1u)) {
+ for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
{
- uint inner = 0u;
- for(; !(!((inner < ColPerThread))); inner = (inner + 1u)) {
+ for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
BCached[inner] = mm_Bsub[k][(tileCol + inner)];
}
}
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
ACached = mm_Asub[(tileRow + innerRow)][k];
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint index = ((innerRow * ColPerThread) + innerCol);
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
}
@@ -158,11 +156,9 @@
}
}
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint index = ((innerRow * ColPerThread) + innerCol);
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
}
diff --git a/test/bug/tint/914.wgsl.expected.msl b/test/bug/tint/914.wgsl.expected.msl
index d9a4b82..eb27d80 100644
--- a/test/bug/tint/914.wgsl.expected.msl
+++ b/test/bug/tint/914.wgsl.expected.msl
@@ -68,155 +68,47 @@
tint_array_wrapper_2 acc = {};
float ACached = 0.0f;
tint_array_wrapper_3 BCached = {};
- {
- uint index = 0u;
- while (true) {
- if (!((index < (RowPerThread * ColPerThread)))) {
- break;
- }
- acc.arr[index] = 0.0f;
- {
- index = (index + 1u);
- }
- }
+ for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
+ acc.arr[index] = 0.0f;
}
uint const ColPerThreadA = (TileInner / 16u);
uint const tileColA = (local_id.x * ColPerThreadA);
uint const RowPerThreadB = (TileInner / 16u);
uint const tileRowB = (local_id.y * RowPerThreadB);
- {
- uint t = 0u;
- while (true) {
- if (!((t < numTiles))) {
- break;
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThreadA))) {
- break;
- }
- uint const inputRow = (tileRow + innerRow);
- uint const inputCol = (tileColA + innerCol);
- tint_symbol_4.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
- {
- innerCol = (innerCol + 1u);
- }
- }
- }
- {
- innerRow = (innerRow + 1u);
- }
- }
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThreadB))) {
- break;
- }
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- uint const inputRow = (tileRowB + innerRow);
- uint const inputCol = (tileCol + innerCol);
- tint_symbol_5.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
- {
- innerCol = (innerCol + 1u);
- }
- }
- }
- {
- innerRow = (innerRow + 1u);
- }
- }
- }
- threadgroup_barrier(mem_flags::mem_threadgroup);
- {
- uint k = 0u;
- while (true) {
- if (!((k < TileInner))) {
- break;
- }
- {
- uint inner = 0u;
- while (true) {
- if (!((inner < ColPerThread))) {
- break;
- }
- BCached.arr[inner] = tint_symbol_5.arr[k].arr[(tileCol + inner)];
- {
- inner = (inner + 1u);
- }
- }
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- uint const index = ((innerRow * ColPerThread) + innerCol);
- acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
- {
- innerCol = (innerCol + 1u);
- }
- }
- }
- {
- innerRow = (innerRow + 1u);
- }
- }
- }
- {
- k = (k + 1u);
- }
- }
- }
- threadgroup_barrier(mem_flags::mem_threadgroup);
- {
- t = (t + 1u);
+ for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+ uint const inputRow = (tileRow + innerRow);
+ uint const inputCol = (tileColA + innerCol);
+ tint_symbol_4.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
}
}
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThread))) {
- break;
+ for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ uint const inputRow = (tileRowB + innerRow);
+ uint const inputCol = (tileCol + innerCol);
+ tint_symbol_5.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
}
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThread))) {
- break;
- }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
+ for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
+ BCached.arr[inner] = tint_symbol_5.arr[k].arr[(tileCol + inner)];
+ }
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * ColPerThread) + innerCol);
- mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
- {
- innerCol = (innerCol + 1u);
- }
+ acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
}
}
- {
- innerRow = (innerRow + 1u);
- }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ }
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ uint const index = ((innerRow * ColPerThread) + innerCol);
+ mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
}
}
return;
diff --git a/test/bug/tint/914.wgsl.expected.wgsl b/test/bug/tint/914.wgsl.expected.wgsl
index ed45fab..17ca3f3 100644
--- a/test/bug/tint/914.wgsl.expected.wgsl
+++ b/test/bug/tint/914.wgsl.expected.wgsl
@@ -65,167 +65,47 @@
var acc : array<f32, 16>;
var ACached : f32;
var BCached : array<f32, 4>;
- {
- var index : u32 = 0u;
- loop {
- if (!((index < (RowPerThread * ColPerThread)))) {
- break;
- }
- acc[index] = 0.0;
-
- continuing {
- index = (index + 1u);
- }
- }
+ for(var index : u32 = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
+ acc[index] = 0.0;
}
let ColPerThreadA : u32 = (TileInner / 16u);
let tileColA : u32 = (local_id.x * ColPerThreadA);
let RowPerThreadB : u32 = (TileInner / 16u);
let tileRowB : u32 = (local_id.y * RowPerThreadB);
- {
- var t : u32 = 0u;
- loop {
- if (!((t < numTiles))) {
- break;
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThreadA))) {
- break;
- }
- let inputRow : u32 = (tileRow + innerRow);
- let inputCol : u32 = (tileColA + innerCol);
- mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
-
- continuing {
- innerCol = (innerCol + 1u);
- }
- }
- }
-
- continuing {
- innerRow = (innerRow + 1u);
- }
- }
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThreadB))) {
- break;
- }
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- let inputRow : u32 = (tileRowB + innerRow);
- let inputCol : u32 = (tileCol + innerCol);
- mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
-
- continuing {
- innerCol = (innerCol + 1u);
- }
- }
- }
-
- continuing {
- innerRow = (innerRow + 1u);
- }
- }
- }
- workgroupBarrier();
- {
- var k : u32 = 0u;
- loop {
- if (!((k < TileInner))) {
- break;
- }
- {
- var inner : u32 = 0u;
- loop {
- if (!((inner < ColPerThread))) {
- break;
- }
- BCached[inner] = mm_Bsub[k][(tileCol + inner)];
-
- continuing {
- inner = (inner + 1u);
- }
- }
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- ACached = mm_Asub[(tileRow + innerRow)][k];
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- let index : u32 = ((innerRow * ColPerThread) + innerCol);
- acc[index] = (acc[index] + (ACached * BCached[innerCol]));
-
- continuing {
- innerCol = (innerCol + 1u);
- }
- }
- }
-
- continuing {
- innerRow = (innerRow + 1u);
- }
- }
- }
-
- continuing {
- k = (k + 1u);
- }
- }
- }
- workgroupBarrier();
-
- continuing {
- t = (t + 1u);
+ for(var t : u32 = 0u; (t < numTiles); t = (t + 1u)) {
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+ let inputRow : u32 = (tileRow + innerRow);
+ let inputCol : u32 = (tileColA + innerCol);
+ mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
}
}
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThread))) {
- break;
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ let inputRow : u32 = (tileRowB + innerRow);
+ let inputCol : u32 = (tileCol + innerCol);
+ mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
}
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThread))) {
- break;
- }
+ }
+ workgroupBarrier();
+ for(var k : u32 = 0u; (k < TileInner); k = (k + 1u)) {
+ for(var inner : u32 = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
+ BCached[inner] = mm_Bsub[k][(tileCol + inner)];
+ }
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ ACached = mm_Asub[(tileRow + innerRow)][k];
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
let index : u32 = ((innerRow * ColPerThread) + innerCol);
- mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
-
- continuing {
- innerCol = (innerCol + 1u);
- }
+ acc[index] = (acc[index] + (ACached * BCached[innerCol]));
}
}
-
- continuing {
- innerRow = (innerRow + 1u);
- }
+ }
+ workgroupBarrier();
+ }
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ let index : u32 = ((innerRow * ColPerThread) + innerCol);
+ mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
}
}
}
diff --git a/test/bug/tint/942.wgsl.expected.hlsl b/test/bug/tint/942.wgsl.expected.hlsl
index ff2739f..695e8ae 100644
--- a/test/bug/tint/942.wgsl.expected.hlsl
+++ b/test/bug/tint/942.wgsl.expected.hlsl
@@ -22,9 +22,13 @@
const uint3 LocalInvocationID = tint_symbol.LocalInvocationID;
const uint local_invocation_index = tint_symbol.local_invocation_index;
if ((local_invocation_index == 0u)) {
- for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
- for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
- tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+ {
+ for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
+ {
+ for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
+ tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+ }
+ }
}
}
}
@@ -37,11 +41,9 @@
const uint scalar_offset_1 = (4u) / 4;
const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[scalar_offset_1 / 4][scalar_offset_1 % 4], 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
{
- uint r = 0u;
- for(; !(!((r < 4u))); r = (r + 1u)) {
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
- uint c = 0u;
- for(; !(!((c < 4u))); c = (c + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
const uint scalar_offset_2 = (0u) / 4;
if ((flip[scalar_offset_2 / 4][scalar_offset_2 % 4] != 0u)) {
@@ -54,11 +56,9 @@
}
GroupMemoryBarrierWithGroupSync();
{
- uint r = 0u;
- for(; !(!((r < 4u))); r = (r + 1u)) {
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
- uint c = 0u;
- for(; !(!((c < 4u))); c = (c + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
const uint scalar_offset_3 = (0u) / 4;
if ((flip[scalar_offset_3 / 4][scalar_offset_3 % 4] != 0u)) {
@@ -79,7 +79,7 @@
uint f = 0u;
while (true) {
const uint scalar_offset_4 = (0u) / 4;
- if (!(!(!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))))) { break; }
+ if (!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))) { break; }
uint i = ((center + f) - filterOffset);
const uint scalar_offset_5 = (0u) / 4;
acc = (acc + ((1.0f / float(params[scalar_offset_5 / 4][scalar_offset_5 % 4])) * tile[r][i]));
diff --git a/test/bug/tint/942.wgsl.expected.msl b/test/bug/tint/942.wgsl.expected.msl
index 16e440d..cecfa13 100644
--- a/test/bug/tint/942.wgsl.expected.msl
+++ b/test/bug/tint/942.wgsl.expected.msl
@@ -25,75 +25,30 @@
uint const filterOffset = ((params.filterDim - 1u) / 2u);
int2 const dims = int2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(0));
int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
- {
- uint r = 0u;
- while (true) {
- if (!((r < 4u))) {
- break;
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+ int2 loadIndex = (baseIndex + int2(int(c), int(r)));
+ if ((flip.value != 0u)) {
+ loadIndex = loadIndex.yx;
}
- {
- uint c = 0u;
- while (true) {
- if (!((c < 4u))) {
- break;
- }
- int2 loadIndex = (baseIndex + int2(int(c), int(r)));
- if ((flip.value != 0u)) {
- loadIndex = loadIndex.yx;
- }
- tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
- {
- c = (c + 1u);
- }
- }
- }
- {
- r = (r + 1u);
- }
+ tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- {
- uint r = 0u;
- while (true) {
- if (!((r < 4u))) {
- break;
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+ int2 writeIndex = (baseIndex + int2(int(c), int(r)));
+ if ((flip.value != 0u)) {
+ writeIndex = writeIndex.yx;
}
- {
- uint c = 0u;
- while (true) {
- if (!((c < 4u))) {
- break;
- }
- int2 writeIndex = (baseIndex + int2(int(c), int(r)));
- if ((flip.value != 0u)) {
- writeIndex = writeIndex.yx;
- }
- uint const center = ((4u * LocalInvocationID.x) + c);
- if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
- float3 acc = float3(0.0f, 0.0f, 0.0f);
- {
- uint f = 0u;
- while (true) {
- if (!((f < params.filterDim))) {
- break;
- }
- uint i = ((center + f) - filterOffset);
- acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
- {
- f = (f + 1u);
- }
- }
- }
- tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
- }
- {
- c = (c + 1u);
- }
+ uint const center = ((4u * LocalInvocationID.x) + c);
+ if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+ float3 acc = float3(0.0f, 0.0f, 0.0f);
+ for(uint f = 0u; (f < params.filterDim); f = (f + 1u)) {
+ uint i = ((center + f) - filterOffset);
+ acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
}
- }
- {
- r = (r + 1u);
+ tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
}
}
}
diff --git a/test/bug/tint/942.wgsl.expected.wgsl b/test/bug/tint/942.wgsl.expected.wgsl
index d261497..b936d4e 100644
--- a/test/bug/tint/942.wgsl.expected.wgsl
+++ b/test/bug/tint/942.wgsl.expected.wgsl
@@ -26,80 +26,30 @@
let filterOffset : u32 = ((params.filterDim - 1u) / 2u);
let dims : vec2<i32> = textureDimensions(inputTex, 0);
let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
- {
- var r : u32 = 0u;
- loop {
- if (!((r < 4u))) {
- break;
+ for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
+ for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
+ var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+ if ((flip.value != 0u)) {
+ loadIndex = loadIndex.yx;
}
- {
- var c : u32 = 0u;
- loop {
- if (!((c < 4u))) {
- break;
- }
- var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
- if ((flip.value != 0u)) {
- loadIndex = loadIndex.yx;
- }
- tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
-
- continuing {
- c = (c + 1u);
- }
- }
- }
-
- continuing {
- r = (r + 1u);
- }
+ tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
}
}
workgroupBarrier();
- {
- var r : u32 = 0u;
- loop {
- if (!((r < 4u))) {
- break;
+ for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
+ for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
+ var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+ if ((flip.value != 0u)) {
+ writeIndex = writeIndex.yx;
}
- {
- var c : u32 = 0u;
- loop {
- if (!((c < 4u))) {
- break;
- }
- var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
- if ((flip.value != 0u)) {
- writeIndex = writeIndex.yx;
- }
- let center : u32 = ((4u * LocalInvocationID.x) + c);
- if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
- var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
- {
- var f : u32 = 0u;
- loop {
- if (!((f < params.filterDim))) {
- break;
- }
- var i : u32 = ((center + f) - filterOffset);
- acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
-
- continuing {
- f = (f + 1u);
- }
- }
- }
- textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
- }
-
- continuing {
- c = (c + 1u);
- }
+ let center : u32 = ((4u * LocalInvocationID.x) + c);
+ if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+ var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
+ for(var f : u32 = 0u; (f < params.filterDim); f = (f + 1u)) {
+ var i : u32 = ((center + f) - filterOffset);
+ acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
}
- }
-
- continuing {
- r = (r + 1u);
+ textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
}
}
}
diff --git a/test/bug/tint/943.spvasm.expected.hlsl b/test/bug/tint/943.spvasm.expected.hlsl
index 47f46cd..a317e46 100644
--- a/test/bug/tint/943.spvasm.expected.hlsl
+++ b/test/bug/tint/943.spvasm.expected.hlsl
@@ -191,10 +191,14 @@
const int x_152 = dimInner;
numTiles = (((x_152 - 1) / 64) + 1);
innerRow = 0;
- for(; (innerRow < 1); innerRow = (innerRow + 1)) {
- innerCol = 0;
- for(; (innerCol < 1); innerCol = (innerCol + 1)) {
- acc[innerRow][innerCol] = 0.0f;
+ {
+ for(; (innerRow < 1); innerRow = (innerRow + 1)) {
+ innerCol = 0;
+ {
+ for(; (innerCol < 1); innerCol = (innerCol + 1)) {
+ acc[innerRow][innerCol] = 0.0f;
+ }
+ }
}
}
const uint x_187 = gl_LocalInvocationID.x;
@@ -202,100 +206,120 @@
const uint x_192 = gl_LocalInvocationID.y;
tileRowB = (asint(x_192) * 1);
t = 0;
- for(; (t < numTiles); t = (t + 1)) {
- innerRow_1 = 0;
- for(; (innerRow_1 < 1); innerRow_1 = (innerRow_1 + 1)) {
- innerCol_1 = 0;
- for(; (innerCol_1 < 64); innerCol_1 = (innerCol_1 + 1)) {
- inputRow = (tileRow + innerRow_1);
- inputCol = (tileColA + innerCol_1);
- const int x_233 = inputRow;
- const int x_234 = inputCol;
- const int x_238 = t;
- const int x_240 = inputCol;
- param_3 = (globalRow + innerRow_1);
- param_4 = ((x_238 * 64) + x_240);
- const float x_244 = mm_readA_i1_i1_(param_3, param_4);
- mm_Asub[x_233][x_234] = x_244;
- }
- }
- innerRow_2 = 0;
- for(; (innerRow_2 < 1); innerRow_2 = (innerRow_2 + 1)) {
- innerCol_2 = 0;
- for(; (innerCol_2 < 1); innerCol_2 = (innerCol_2 + 1)) {
- inputRow_1 = (tileRowB + innerRow_2);
- inputCol_1 = (tileCol + innerCol_2);
- const int x_278 = inputRow_1;
- const int x_279 = inputCol_1;
- const int x_284 = globalCol;
- const int x_285 = innerCol_2;
- param_5 = ((t * 64) + inputRow_1);
- param_6 = (x_284 + x_285);
- const float x_289 = mm_readB_i1_i1_(param_5, param_6);
- mm_Bsub[x_278][x_279] = x_289;
- }
- }
- GroupMemoryBarrierWithGroupSync();
- k = 0;
- for(; (k < 64); k = (k + 1)) {
- inner = 0;
- for(; (inner < 1); inner = (inner + 1)) {
- const int x_314 = inner;
- const float x_320 = mm_Bsub[k][(tileCol + inner)];
- BCached[x_314] = x_320;
- }
- innerRow_3 = 0;
- for(; (innerRow_3 < 1); innerRow_3 = (innerRow_3 + 1)) {
- const float x_338 = mm_Asub[(tileRow + innerRow_3)][k];
- ACached = x_338;
- innerCol_3 = 0;
- for(; (innerCol_3 < 1); innerCol_3 = (innerCol_3 + 1)) {
- const int x_347 = innerRow_3;
- const int x_348 = innerCol_3;
- const float x_349 = ACached;
- const float x_352 = BCached[innerCol_3];
- const float x_355 = acc[x_347][x_348];
- acc[x_347][x_348] = (x_355 + (x_349 * x_352));
+ {
+ for(; (t < numTiles); t = (t + 1)) {
+ innerRow_1 = 0;
+ {
+ for(; (innerRow_1 < 1); innerRow_1 = (innerRow_1 + 1)) {
+ innerCol_1 = 0;
+ {
+ for(; (innerCol_1 < 64); innerCol_1 = (innerCol_1 + 1)) {
+ inputRow = (tileRow + innerRow_1);
+ inputCol = (tileColA + innerCol_1);
+ const int x_233 = inputRow;
+ const int x_234 = inputCol;
+ const int x_238 = t;
+ const int x_240 = inputCol;
+ param_3 = (globalRow + innerRow_1);
+ param_4 = ((x_238 * 64) + x_240);
+ const float x_244 = mm_readA_i1_i1_(param_3, param_4);
+ mm_Asub[x_233][x_234] = x_244;
+ }
+ }
}
}
+ innerRow_2 = 0;
+ {
+ for(; (innerRow_2 < 1); innerRow_2 = (innerRow_2 + 1)) {
+ innerCol_2 = 0;
+ {
+ for(; (innerCol_2 < 1); innerCol_2 = (innerCol_2 + 1)) {
+ inputRow_1 = (tileRowB + innerRow_2);
+ inputCol_1 = (tileCol + innerCol_2);
+ const int x_278 = inputRow_1;
+ const int x_279 = inputCol_1;
+ const int x_284 = globalCol;
+ const int x_285 = innerCol_2;
+ param_5 = ((t * 64) + inputRow_1);
+ param_6 = (x_284 + x_285);
+ const float x_289 = mm_readB_i1_i1_(param_5, param_6);
+ mm_Bsub[x_278][x_279] = x_289;
+ }
+ }
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ k = 0;
+ {
+ for(; (k < 64); k = (k + 1)) {
+ inner = 0;
+ {
+ for(; (inner < 1); inner = (inner + 1)) {
+ const int x_314 = inner;
+ const float x_320 = mm_Bsub[k][(tileCol + inner)];
+ BCached[x_314] = x_320;
+ }
+ }
+ innerRow_3 = 0;
+ {
+ for(; (innerRow_3 < 1); innerRow_3 = (innerRow_3 + 1)) {
+ const float x_338 = mm_Asub[(tileRow + innerRow_3)][k];
+ ACached = x_338;
+ innerCol_3 = 0;
+ {
+ for(; (innerCol_3 < 1); innerCol_3 = (innerCol_3 + 1)) {
+ const int x_347 = innerRow_3;
+ const int x_348 = innerCol_3;
+ const float x_349 = ACached;
+ const float x_352 = BCached[innerCol_3];
+ const float x_355 = acc[x_347][x_348];
+ acc[x_347][x_348] = (x_355 + (x_349 * x_352));
+ }
+ }
+ }
+ }
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
}
- GroupMemoryBarrierWithGroupSync();
}
innerRow_4 = 0;
- for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
- innerCol_4 = 0;
- while (true) {
- bool x_393 = false;
- bool x_394_phi = false;
- if ((innerCol_4 < 1)) {
- } else {
- break;
- }
- const int x_382 = globalCol;
- const int x_383 = innerCol_4;
- const int x_385 = dimBOuter;
- const bool x_386 = ((x_382 + x_383) < x_385);
- x_394_phi = x_386;
- if (x_386) {
- const int x_389 = globalRow;
- const int x_390 = innerRow_4;
- const int x_392 = dimAOuter;
- x_393 = ((x_389 + x_390) < x_392);
- x_394_phi = x_393;
- }
- if (x_394_phi) {
- const int x_400 = globalCol;
- const int x_401 = innerCol_4;
- const int x_403 = innerRow_4;
- const int x_404 = innerCol_4;
- param_7 = (globalRow + innerRow_4);
- param_8 = (x_400 + x_401);
- const float x_409 = acc[x_403][x_404];
- param_9 = x_409;
- mm_write_i1_i1_f1_(param_7, param_8, param_9);
- }
- {
- innerCol_4 = (innerCol_4 + 1);
+ {
+ for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
+ innerCol_4 = 0;
+ while (true) {
+ bool x_393 = false;
+ bool x_394_phi = false;
+ if ((innerCol_4 < 1)) {
+ } else {
+ break;
+ }
+ const int x_382 = globalCol;
+ const int x_383 = innerCol_4;
+ const int x_385 = dimBOuter;
+ const bool x_386 = ((x_382 + x_383) < x_385);
+ x_394_phi = x_386;
+ if (x_386) {
+ const int x_389 = globalRow;
+ const int x_390 = innerRow_4;
+ const int x_392 = dimAOuter;
+ x_393 = ((x_389 + x_390) < x_392);
+ x_394_phi = x_393;
+ }
+ if (x_394_phi) {
+ const int x_400 = globalCol;
+ const int x_401 = innerCol_4;
+ const int x_403 = innerRow_4;
+ const int x_404 = innerCol_4;
+ param_7 = (globalRow + innerRow_4);
+ param_8 = (x_400 + x_401);
+ const float x_409 = acc[x_403][x_404];
+ param_9 = x_409;
+ mm_write_i1_i1_f1_(param_7, param_8, param_9);
+ }
+ {
+ innerCol_4 = (innerCol_4 + 1);
+ }
}
}
}
@@ -336,14 +360,20 @@
const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
const uint local_invocation_index = tint_symbol.local_invocation_index;
if ((local_invocation_index == 0u)) {
- for(int i = 0; (i < 64); i = (i + 1)) {
- for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
- mm_Asub[i][i_1] = 0.0f;
+ {
+ for(int i = 0; (i < 64); i = (i + 1)) {
+ {
+ for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+ mm_Asub[i][i_1] = 0.0f;
+ }
+ }
}
}
- for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
- const float tint_symbol_6[1] = (float[1])0;
- mm_Bsub[i_2] = tint_symbol_6;
+ {
+ for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+ const float tint_symbol_6[1] = (float[1])0;
+ mm_Bsub[i_2] = tint_symbol_6;
+ }
}
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/bug/tint/948.wgsl.expected.hlsl b/test/bug/tint/948.wgsl.expected.hlsl
index 2454c3b..d6f3a02 100644
--- a/test/bug/tint/948.wgsl.expected.hlsl
+++ b/test/bug/tint/948.wgsl.expected.hlsl
@@ -67,89 +67,93 @@
const float2 x_111 = asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
stageUnits = (float2(1.0f, 1.0f) / x_111);
i = 0;
- for(; (i < 2); i = (i + 1)) {
- switch(i) {
- case 1: {
- const float2 x_150 = tileID;
- const uint scalar_offset_4 = (88u) / 4;
- uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
- const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
- const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
- frameID_1 = x_156.x;
- break;
- }
- case 0: {
- const float2 x_136 = tileID;
- const uint scalar_offset_5 = (88u) / 4;
- uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
- const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
- const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
- frameID_1 = x_142.x;
- break;
- }
- default: {
- break;
- }
- }
- const float x_166 = frameID_1;
- const uint scalar_offset_6 = (108u) / 4;
- const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
- const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
- animationData = x_172;
- const float x_174 = animationData.y;
- if ((x_174 > 0.0f)) {
- const uint scalar_offset_7 = (0u) / 4;
- const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
- const float x_184 = animationData.z;
- mt = ((x_181 * x_184) % 1.0f);
- f = 0.0f;
- for(; (f < 8.0f); f = (f + 1.0f)) {
- const float x_197 = animationData.y;
- if ((x_197 > mt)) {
- const float x_203 = animationData.x;
- frameID_1 = x_203;
+ {
+ for(; (i < 2); i = (i + 1)) {
+ switch(i) {
+ case 1: {
+ const float2 x_150 = tileID;
+ const uint scalar_offset_4 = (88u) / 4;
+ uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
+ const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
+ const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
+ frameID_1 = x_156.x;
break;
}
- const float x_208 = frameID_1;
- const uint scalar_offset_8 = (108u) / 4;
- const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
- const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
- animationData = x_217;
+ case 0: {
+ const float2 x_136 = tileID;
+ const uint scalar_offset_5 = (88u) / 4;
+ uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
+ const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
+ const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
+ frameID_1 = x_142.x;
+ break;
+ }
+ default: {
+ break;
+ }
}
- }
- param = (frameID_1 + 0.5f);
- const float4x4 x_225 = getFrameData_f1_(param);
- frameData = x_225;
- const float4 x_228 = frameData[0];
- const uint scalar_offset_9 = (96u) / 4;
- uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
- const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
- frameSize = (float2(x_228.w, x_228.z) / x_231);
- const float4 x_235 = frameData[0];
- offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
- const float4 x_241 = frameData[2];
- const float4 x_244 = frameData[0];
- ratio = (float2(x_241.x, x_241.y) / float2(x_244.w, x_244.z));
- const float x_248 = frameData[2].z;
- if ((x_248 == 1.0f)) {
- const float2 x_252 = tileUV;
- tileUV = float2(x_252.y, x_252.x);
- }
- if ((i == 0)) {
- const float4 x_268 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
- color = x_268;
- } else {
- const float4 x_279 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
- nc = x_279;
- const float x_283 = color.w;
- const float x_285 = nc.w;
- alpha = min((x_283 + x_285), 1.0f);
- const float4 x_290 = color;
- const float4 x_292 = nc;
- const float x_295 = nc.w;
- mixed = lerp(float3(x_290.x, x_290.y, x_290.z), float3(x_292.x, x_292.y, x_292.z), float3(x_295, x_295, x_295));
- const float3 x_298 = mixed;
- color = float4(x_298.x, x_298.y, x_298.z, alpha);
+ const float x_166 = frameID_1;
+ const uint scalar_offset_6 = (108u) / 4;
+ const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+ const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
+ animationData = x_172;
+ const float x_174 = animationData.y;
+ if ((x_174 > 0.0f)) {
+ const uint scalar_offset_7 = (0u) / 4;
+ const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ const float x_184 = animationData.z;
+ mt = ((x_181 * x_184) % 1.0f);
+ f = 0.0f;
+ {
+ for(; (f < 8.0f); f = (f + 1.0f)) {
+ const float x_197 = animationData.y;
+ if ((x_197 > mt)) {
+ const float x_203 = animationData.x;
+ frameID_1 = x_203;
+ break;
+ }
+ const float x_208 = frameID_1;
+ const uint scalar_offset_8 = (108u) / 4;
+ const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+ const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
+ animationData = x_217;
+ }
+ }
+ }
+ param = (frameID_1 + 0.5f);
+ const float4x4 x_225 = getFrameData_f1_(param);
+ frameData = x_225;
+ const float4 x_228 = frameData[0];
+ const uint scalar_offset_9 = (96u) / 4;
+ uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
+ const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
+ frameSize = (float2(x_228.w, x_228.z) / x_231);
+ const float4 x_235 = frameData[0];
+ offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
+ const float4 x_241 = frameData[2];
+ const float4 x_244 = frameData[0];
+ ratio = (float2(x_241.x, x_241.y) / float2(x_244.w, x_244.z));
+ const float x_248 = frameData[2].z;
+ if ((x_248 == 1.0f)) {
+ const float2 x_252 = tileUV;
+ tileUV = float2(x_252.y, x_252.x);
+ }
+ if ((i == 0)) {
+ const float4 x_268 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
+ color = x_268;
+ } else {
+ const float4 x_279 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
+ nc = x_279;
+ const float x_283 = color.w;
+ const float x_285 = nc.w;
+ alpha = min((x_283 + x_285), 1.0f);
+ const float4 x_290 = color;
+ const float4 x_292 = nc;
+ const float x_295 = nc.w;
+ mixed = lerp(float3(x_290.x, x_290.y, x_290.z), float3(x_292.x, x_292.y, x_292.z), float3(x_295, x_295, x_295));
+ const float3 x_298 = mixed;
+ color = float4(x_298.x, x_298.y, x_298.z, alpha);
+ }
}
}
const uint scalar_offset_10 = (112u) / 4;
diff --git a/test/bug/tint/949.wgsl.expected.hlsl b/test/bug/tint/949.wgsl.expected.hlsl
index 926e020..2bb720a 100644
--- a/test/bug/tint/949.wgsl.expected.hlsl
+++ b/test/bug/tint/949.wgsl.expected.hlsl
@@ -249,20 +249,22 @@
lastSampledHeight = 1.0f;
currSampledHeight = 1.0f;
i = 0;
- for(; (i < 15); i = (i + 1)) {
- const float4 x_397 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + vCurrOffset));
- currSampledHeight = x_397.w;
- if ((currSampledHeight > currRayHeight)) {
- delta1 = (currSampledHeight - currRayHeight);
- delta2 = ((currRayHeight + stepSize) - lastSampledHeight);
- ratio = (delta1 / (delta1 + delta2));
- vCurrOffset = ((vLastOffset * ratio) + (vCurrOffset * (1.0f - ratio)));
- break;
- } else {
- currRayHeight = (currRayHeight - stepSize);
- vLastOffset = vCurrOffset;
- vCurrOffset = (vCurrOffset + (vMaxOffset * stepSize));
- lastSampledHeight = currSampledHeight;
+ {
+ for(; (i < 15); i = (i + 1)) {
+ const float4 x_397 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + vCurrOffset));
+ currSampledHeight = x_397.w;
+ if ((currSampledHeight > currRayHeight)) {
+ delta1 = (currSampledHeight - currRayHeight);
+ delta2 = ((currRayHeight + stepSize) - lastSampledHeight);
+ ratio = (delta1 / (delta1 + delta2));
+ vCurrOffset = ((vLastOffset * ratio) + (vCurrOffset * (1.0f - ratio)));
+ break;
+ } else {
+ currRayHeight = (currRayHeight - stepSize);
+ vLastOffset = vCurrOffset;
+ vCurrOffset = (vCurrOffset + (vMaxOffset * stepSize));
+ lastSampledHeight = currSampledHeight;
+ }
}
}
parallaxOcclusion_0 = vCurrOffset;
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
index 6496cf5..a222773 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
@@ -32,8 +32,7 @@
void foo() {
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_int3(v3i, i, 1);
Set_uint4(v4u, i, 1u);
@@ -45,8 +44,7 @@
[numthreads(1, 1, 1)]
void main() {
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
index ab45d31..2ab266f 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
@@ -2,20 +2,11 @@
using namespace metal;
void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) {
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- (*(tint_symbol_1))[i] = 1.0f;
- (*(tint_symbol_2))[i] = 1;
- (*(tint_symbol_3))[i] = 1u;
- (*(tint_symbol_4))[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ (*(tint_symbol_1))[i] = 1.0f;
+ (*(tint_symbol_2))[i] = 1;
+ (*(tint_symbol_3))[i] = 1u;
+ (*(tint_symbol_4))[i] = true;
}
}
@@ -24,17 +15,8 @@
thread int3 tint_symbol_6 = 0;
thread uint4 tint_symbol_7 = 0u;
thread bool2 tint_symbol_8 = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
index bc302f1..def10cb 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
@@ -7,37 +7,17 @@
var<private> v2b : vec2<bool>;
fn foo() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
}
}
[[stage(compute), workgroup_size(1, 1, 1)]]
fn main() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- foo();
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
index 33a99b9..ea05754 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
@@ -14,8 +14,7 @@
[numthreads(1, 1, 1)]
void main() {
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
index 2685e75..44a4fa0 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
@@ -14,17 +14,8 @@
thread int3 tint_symbol_6 = 0;
thread uint4 tint_symbol_7 = 0u;
thread bool2 tint_symbol_8 = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
index ae1667c..21db9dd 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
@@ -16,17 +16,7 @@
[[stage(compute), workgroup_size(1, 1, 1)]]
fn main() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- foo();
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
index c8fa518..85ffc25 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
@@ -97,8 +97,7 @@
bool3 v3b = bool3(false, false, false);
bool4 v4b = bool4(false, false, false, false);
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_float3(v3f, i, 1.0f);
Set_float4(v4f, i, 1.0f);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
index a460250..b622316 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
@@ -14,28 +14,19 @@
bool2 v2b = false;
bool3 v3b = false;
bool4 v4b = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0f;
- v3f[i] = 1.0f;
- v4f[i] = 1.0f;
- v2i[i] = 1;
- v3i[i] = 1;
- v4i[i] = 1;
- v2u[i] = 1u;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v2b[i] = true;
- v3b[i] = true;
- v4b[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0f;
+ v3f[i] = 1.0f;
+ v4f[i] = 1.0f;
+ v2i[i] = 1;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v2u[i] = 1u;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v3b[i] = true;
+ v4b[i] = true;
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
index 7d73628..9f8f2e0 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
@@ -12,28 +12,18 @@
var v2b : vec2<bool>;
var v3b : vec3<bool>;
var v4b : vec4<bool>;
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v3f[i] = 1.0;
- v4f[i] = 1.0;
- v2i[i] = 1;
- v3i[i] = 1;
- v4i[i] = 1;
- v2u[i] = 1u;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v2b[i] = true;
- v3b[i] = true;
- v4b[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v3f[i] = 1.0;
+ v4f[i] = 1.0;
+ v2i[i] = 1;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v2u[i] = 1u;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v3b[i] = true;
+ v4b[i] = true;
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
index cbc4e6a..9c60097 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
@@ -36,8 +36,7 @@
bool2 v2b = bool2(false, false);
bool2 v2b_2 = bool2(false, false);
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_int3(v3i, i, 1);
Set_uint4(v4u, i, 1u);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
index 9f28e5c..4c124ca 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
@@ -10,24 +10,15 @@
uint4 v4u_2 = 0u;
bool2 v2b = false;
bool2 v2b_2 = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0f;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
- v2f_2[i] = 1.0f;
- v3i_2[i] = 1;
- v4u_2[i] = 1u;
- v2b_2[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0f;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v2f_2[i] = 1.0f;
+ v3i_2[i] = 1;
+ v4u_2[i] = 1u;
+ v2b_2[i] = true;
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
index 0229323..74950ca 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
@@ -8,24 +8,14 @@
var v4u_2 : vec4<u32>;
var v2b : vec2<bool>;
var v2b_2 : vec2<bool>;
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
- v2f_2[i] = 1.0;
- v3i_2[i] = 1;
- v4u_2[i] = 1u;
- v2b_2[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v2f_2[i] = 1.0;
+ v3i_2[i] = 1;
+ v4u_2[i] = 1u;
+ v2b_2[i] = true;
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
index aed923f..2e6de37 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
@@ -37,8 +37,7 @@
bool3 v3b = bool3(false, false, false);
bool4 v4b = bool4(false, false, false, false);
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_int2(v2i, i, 1);
Set_uint2(v2u, i, 1u);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
index fb53199..2113ef6 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
@@ -14,20 +14,11 @@
bool2 v2b = false;
bool3 v3b = false;
bool4 v4b = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0f;
- v2i[i] = 1;
- v2u[i] = 1u;
- v2b[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0f;
+ v2i[i] = 1;
+ v2u[i] = 1u;
+ v2b[i] = true;
}
int i = 0;
v3f[i] = 1.0f;
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
index 8c54146..1fada75 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
@@ -12,21 +12,11 @@
var v2b : vec2<bool>;
var v3b : vec3<bool>;
var v4b : vec4<bool>;
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v2i[i] = 1;
- v2u[i] = 1u;
- v2b[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v2i[i] = 1;
+ v2u[i] = 1u;
+ v2b[i] = true;
}
var i : i32 = 0;
v3f[i] = 1.0;
diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl
index d98a316..4643faa 100644
--- a/test/samples/compute_boids.wgsl.expected.hlsl
+++ b/test/samples/compute_boids.wgsl.expected.hlsl
@@ -53,8 +53,7 @@
float2 pos = float2(0.0f, 0.0f);
float2 vel = float2(0.0f, 0.0f);
{
- uint i = 0u;
- for(; !(!((i < 5u))); i = (i + 1u)) {
+ for(uint i = 0u; (i < 5u); i = (i + 1u)) {
if ((i == index)) {
continue;
}
diff --git a/test/samples/compute_boids.wgsl.expected.msl b/test/samples/compute_boids.wgsl.expected.msl
index dae2f1f..cc9a856 100644
--- a/test/samples/compute_boids.wgsl.expected.msl
+++ b/test/samples/compute_boids.wgsl.expected.msl
@@ -61,34 +61,22 @@
int cVelCount = 0;
float2 pos = 0.0f;
float2 vel = 0.0f;
- {
- uint i = 0u;
- while (true) {
- if (!((i < 5u))) {
- break;
- }
- if ((i == index)) {
- {
- i = (i + 1u);
- }
- continue;
- }
- pos = particlesA.particles.arr[i].pos.xy;
- vel = particlesA.particles.arr[i].vel.xy;
- if ((distance(pos, vPos) < params.rule1Distance)) {
- cMass = (cMass + pos);
- cMassCount = (cMassCount + 1);
- }
- if ((distance(pos, vPos) < params.rule2Distance)) {
- colVel = (colVel - (pos - vPos));
- }
- if ((distance(pos, vPos) < params.rule3Distance)) {
- cVel = (cVel + vel);
- cVelCount = (cVelCount + 1);
- }
- {
- i = (i + 1u);
- }
+ for(uint i = 0u; (i < 5u); i = (i + 1u)) {
+ if ((i == index)) {
+ continue;
+ }
+ pos = particlesA.particles.arr[i].pos.xy;
+ vel = particlesA.particles.arr[i].vel.xy;
+ if ((distance(pos, vPos) < params.rule1Distance)) {
+ cMass = (cMass + pos);
+ cMassCount = (cMassCount + 1);
+ }
+ if ((distance(pos, vPos) < params.rule2Distance)) {
+ colVel = (colVel - (pos - vPos));
+ }
+ if ((distance(pos, vPos) < params.rule3Distance)) {
+ cVel = (cVel + vel);
+ cVelCount = (cVelCount + 1);
}
}
if ((cMassCount > 0)) {
diff --git a/test/samples/compute_boids.wgsl.expected.wgsl b/test/samples/compute_boids.wgsl.expected.wgsl
index 09cb30f..dd981f4 100644
--- a/test/samples/compute_boids.wgsl.expected.wgsl
+++ b/test/samples/compute_boids.wgsl.expected.wgsl
@@ -52,32 +52,22 @@
var cVelCount : i32 = 0;
var pos : vec2<f32>;
var vel : vec2<f32>;
- {
- var i : u32 = 0u;
- loop {
- if (!((i < 5u))) {
- break;
- }
- if ((i == index)) {
- continue;
- }
- pos = particlesA.particles[i].pos.xy;
- vel = particlesA.particles[i].vel.xy;
- if ((distance(pos, vPos) < params.rule1Distance)) {
- cMass = (cMass + pos);
- cMassCount = (cMassCount + 1);
- }
- if ((distance(pos, vPos) < params.rule2Distance)) {
- colVel = (colVel - (pos - vPos));
- }
- if ((distance(pos, vPos) < params.rule3Distance)) {
- cVel = (cVel + vel);
- cVelCount = (cVelCount + 1);
- }
-
- continuing {
- i = (i + 1u);
- }
+ for(var i : u32 = 0u; (i < 5u); i = (i + 1u)) {
+ if ((i == index)) {
+ continue;
+ }
+ pos = particlesA.particles[i].pos.xy;
+ vel = particlesA.particles[i].vel.xy;
+ if ((distance(pos, vPos) < params.rule1Distance)) {
+ cMass = (cMass + pos);
+ cMassCount = (cMassCount + 1);
+ }
+ if ((distance(pos, vPos) < params.rule2Distance)) {
+ colVel = (colVel - (pos - vPos));
+ }
+ if ((distance(pos, vPos) < params.rule3Distance)) {
+ cVel = (cVel + vel);
+ cVelCount = (cVelCount + 1);
}
}
if ((cMassCount > 0)) {
diff --git a/test/statements/for/basic.wgsl.expected.hlsl b/test/statements/for/basic.wgsl.expected.hlsl
index c1e6281..f215f4b 100644
--- a/test/statements/for/basic.wgsl.expected.hlsl
+++ b/test/statements/for/basic.wgsl.expected.hlsl
@@ -8,8 +8,7 @@
void f() {
{
- int i = 0;
- for(; !(!((i < 5))); i = (i + 1)) {
+ for(int i = 0; (i < 5); i = (i + 1)) {
some_loop_body();
}
}
diff --git a/test/statements/for/basic.wgsl.expected.msl b/test/statements/for/basic.wgsl.expected.msl
index d2e498f..e36af9f 100644
--- a/test/statements/for/basic.wgsl.expected.msl
+++ b/test/statements/for/basic.wgsl.expected.msl
@@ -5,17 +5,8 @@
}
void f() {
- {
- int i = 0;
- while (true) {
- if (!((i < 5))) {
- break;
- }
- some_loop_body();
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 5); i = (i + 1)) {
+ some_loop_body();
}
}
diff --git a/test/statements/for/basic.wgsl.expected.wgsl b/test/statements/for/basic.wgsl.expected.wgsl
index 6a66161..3f9a88a 100644
--- a/test/statements/for/basic.wgsl.expected.wgsl
+++ b/test/statements/for/basic.wgsl.expected.wgsl
@@ -2,17 +2,7 @@
}
fn f() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 5))) {
- break;
- }
- some_loop_body();
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 5); i = (i + 1)) {
+ some_loop_body();
}
}
diff --git a/test/statements/for/complex.wgsl.expected.hlsl b/test/statements/for/complex.wgsl.expected.hlsl
index a3197c1..b48727f 100644
--- a/test/statements/for/complex.wgsl.expected.hlsl
+++ b/test/statements/for/complex.wgsl.expected.hlsl
@@ -15,7 +15,7 @@
if (tint_tmp) {
tint_tmp = (j < 10);
}
- if (!(!(!((tint_tmp))))) { break; }
+ if (!((tint_tmp))) { break; }
some_loop_body();
j = (i * 30);
i = (i + 1);
diff --git a/test/statements/for/complex.wgsl.expected.msl b/test/statements/for/complex.wgsl.expected.msl
index 2030618..9ed6331 100644
--- a/test/statements/for/complex.wgsl.expected.msl
+++ b/test/statements/for/complex.wgsl.expected.msl
@@ -6,18 +6,9 @@
void f() {
int j = 0;
- {
- int i = 0;
- while (true) {
- if (!(((i < 5) && (j < 10)))) {
- break;
- }
- some_loop_body();
- j = (i * 30);
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
+ some_loop_body();
+ j = (i * 30);
}
}
diff --git a/test/statements/for/complex.wgsl.expected.wgsl b/test/statements/for/complex.wgsl.expected.wgsl
index 1f3ce92..32bc4e0 100644
--- a/test/statements/for/complex.wgsl.expected.wgsl
+++ b/test/statements/for/complex.wgsl.expected.wgsl
@@ -3,18 +3,8 @@
fn f() {
var j : i32;
- {
- var i : i32 = 0;
- loop {
- if (!(((i < 5) && (j < 10)))) {
- break;
- }
- some_loop_body();
- j = (i * 30);
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
+ some_loop_body();
+ j = (i * 30);
}
}
diff --git a/test/statements/for/condition.wgsl.expected.hlsl b/test/statements/for/condition.wgsl.expected.hlsl
index 5f51493..e4456e0 100644
--- a/test/statements/for/condition.wgsl.expected.hlsl
+++ b/test/statements/for/condition.wgsl.expected.hlsl
@@ -5,6 +5,8 @@
void f() {
int i = 0;
- for(; !(!((i < 4))); ) {
+ {
+ for(; (i < 4); ) {
+ }
}
}
diff --git a/test/statements/for/condition.wgsl.expected.msl b/test/statements/for/condition.wgsl.expected.msl
index c99148a..839541b 100644
--- a/test/statements/for/condition.wgsl.expected.msl
+++ b/test/statements/for/condition.wgsl.expected.msl
@@ -3,10 +3,7 @@
using namespace metal;
void f() {
int i = 0;
- while (true) {
- if (!((i < 4))) {
- break;
- }
+ for(; (i < 4); ) {
}
}
diff --git a/test/statements/for/condition.wgsl.expected.wgsl b/test/statements/for/condition.wgsl.expected.wgsl
index 14f48d5..7fc3c6f 100644
--- a/test/statements/for/condition.wgsl.expected.wgsl
+++ b/test/statements/for/condition.wgsl.expected.wgsl
@@ -1,8 +1,5 @@
fn f() {
var i : i32;
- loop {
- if (!((i < 4))) {
- break;
- }
+ for(; (i < 4); ;) {
}
}
diff --git a/test/statements/for/continuing.wgsl.expected.hlsl b/test/statements/for/continuing.wgsl.expected.hlsl
index 28b8b3b..b3ec736 100644
--- a/test/statements/for/continuing.wgsl.expected.hlsl
+++ b/test/statements/for/continuing.wgsl.expected.hlsl
@@ -5,9 +5,8 @@
void f() {
int i = 0;
- while (true) {
- {
- i = (i + 1);
+ {
+ for(; ; i = (i + 1)) {
}
}
}
diff --git a/test/statements/for/continuing.wgsl.expected.msl b/test/statements/for/continuing.wgsl.expected.msl
index 67a069f..a288b5d 100644
--- a/test/statements/for/continuing.wgsl.expected.msl
+++ b/test/statements/for/continuing.wgsl.expected.msl
@@ -3,10 +3,7 @@
using namespace metal;
void f() {
int i = 0;
- while (true) {
- {
- i = (i + 1);
- }
+ for(; ; i = (i + 1)) {
}
}
diff --git a/test/statements/for/continuing.wgsl.expected.wgsl b/test/statements/for/continuing.wgsl.expected.wgsl
index f3faf03..c76b514 100644
--- a/test/statements/for/continuing.wgsl.expected.wgsl
+++ b/test/statements/for/continuing.wgsl.expected.wgsl
@@ -1,9 +1,5 @@
fn f() {
var i : i32;
- loop {
-
- continuing {
- i = (i + 1);
- }
+ for(; ; i = (i + 1)) {
}
}
diff --git a/test/statements/for/empty.wgsl.expected.hlsl b/test/statements/for/empty.wgsl.expected.hlsl
index b642595..cd85e1a 100644
--- a/test/statements/for/empty.wgsl.expected.hlsl
+++ b/test/statements/for/empty.wgsl.expected.hlsl
@@ -4,6 +4,8 @@
}
void f() {
- while (true) {
+ {
+ for(; ; ) {
+ }
}
}
diff --git a/test/statements/for/empty.wgsl.expected.msl b/test/statements/for/empty.wgsl.expected.msl
index 622427b..e716521 100644
--- a/test/statements/for/empty.wgsl.expected.msl
+++ b/test/statements/for/empty.wgsl.expected.msl
@@ -2,7 +2,7 @@
using namespace metal;
void f() {
- while (true) {
+ for(; ; ) {
}
}
diff --git a/test/statements/for/empty.wgsl.expected.wgsl b/test/statements/for/empty.wgsl.expected.wgsl
index 3674ab1..bb4e48b 100644
--- a/test/statements/for/empty.wgsl.expected.wgsl
+++ b/test/statements/for/empty.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
fn f() {
- loop {
+ for(; ; ;) {
}
}
diff --git a/test/statements/for/initializer.wgsl.expected.hlsl b/test/statements/for/initializer.wgsl.expected.hlsl
index 8311c7b..e614afd 100644
--- a/test/statements/for/initializer.wgsl.expected.hlsl
+++ b/test/statements/for/initializer.wgsl.expected.hlsl
@@ -5,8 +5,7 @@
void f() {
{
- int i = 0;
- while (true) {
+ for(int i = 0; ; ) {
}
}
}
diff --git a/test/statements/for/initializer.wgsl.expected.msl b/test/statements/for/initializer.wgsl.expected.msl
index 70c5111..90082d6 100644
--- a/test/statements/for/initializer.wgsl.expected.msl
+++ b/test/statements/for/initializer.wgsl.expected.msl
@@ -2,10 +2,7 @@
using namespace metal;
void f() {
- {
- int i = 0;
- while (true) {
- }
+ for(int i = 0; ; ) {
}
}
diff --git a/test/statements/for/initializer.wgsl.expected.wgsl b/test/statements/for/initializer.wgsl.expected.wgsl
index e021556..ef9b991 100644
--- a/test/statements/for/initializer.wgsl.expected.wgsl
+++ b/test/statements/for/initializer.wgsl.expected.wgsl
@@ -1,7 +1,4 @@
fn f() {
- {
- var i : i32 = 0;
- loop {
- }
+ for(var i : i32 = 0; ; ;) {
}
}
diff --git a/test/statements/for/scoping.wgsl.expected.hlsl b/test/statements/for/scoping.wgsl.expected.hlsl
index 459eb43..143b1e3 100644
--- a/test/statements/for/scoping.wgsl.expected.hlsl
+++ b/test/statements/for/scoping.wgsl.expected.hlsl
@@ -5,8 +5,7 @@
void f() {
{
- int must_not_collide = 0;
- while (true) {
+ for(int must_not_collide = 0; ; ) {
}
}
int must_not_collide = 0;
diff --git a/test/statements/for/scoping.wgsl.expected.msl b/test/statements/for/scoping.wgsl.expected.msl
index d5c3f3a..0572921 100644
--- a/test/statements/for/scoping.wgsl.expected.msl
+++ b/test/statements/for/scoping.wgsl.expected.msl
@@ -2,10 +2,7 @@
using namespace metal;
void f() {
- {
- int must_not_collide = 0;
- while (true) {
- }
+ for(int must_not_collide = 0; ; ) {
}
int must_not_collide = 0;
}
diff --git a/test/statements/for/scoping.wgsl.expected.wgsl b/test/statements/for/scoping.wgsl.expected.wgsl
index ec202ae..0cac852 100644
--- a/test/statements/for/scoping.wgsl.expected.wgsl
+++ b/test/statements/for/scoping.wgsl.expected.wgsl
@@ -1,8 +1,5 @@
fn f() {
- {
- var must_not_collide : i32 = 0;
- loop {
- }
+ for(var must_not_collide : i32 = 0; ; ;) {
}
var must_not_collide : i32;
}
