encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-12-09 21:47:47 +00:00
Code Issues Packages Projects Releases Wiki Activity

writer/msl: Rework string printing

Browse Source 
Add `out` parameters to expression and type generators.

Use the new helper classes in TextGenerator.
Cleans up bad formatting.

Prepares the writer generating 'pre' statements, required for atomics.

If-else statements are generated slightly differently. This is done so that 'pre' statements for the else conditions are scoped correctly. This is identical to the HLSL writer.

Bug tint:892

Change-Id: I4c6e96c90673ba30898b3682bf3198497d63a2d4
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56067
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton
2021-06-28 15:30:57 +00:00
committed by Tint LUCI CQ
parent a5715e3320
commit f24b37e122
319 changed files with 1126 additions and 1081 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
test/bug/dawn/947.wgsl.expected.msl
View File

@@ -39,8 +39,8 @@ vertex tint_symbol_1 vs_main(uint VertexIndex [[vertex_id]], constant Uniforms&
fragment tint_symbol_4 fs_main(texture2d<float, access::sample> tint_symbol_7 [[texture(2)]], sampler tint_symbol_8 [[sampler(1)]], tint_symbol_3 tint_symbol_2 [[stage_in]]) {
float2 const texcoord = tint_symbol_2.texcoord;
float2 clampedTexcoord = clamp(texcoord, float2(0.0f, 0.0f), float2(1.0f, 1.0f));
if (!( all((clampedTexcoord == texcoord)))) {
float2 clampedTexcoord = clamp(texcoord, float2(0.0f, 0.0f), float2(1.0f, 1.0f));
if (!(all((clampedTexcoord == texcoord)))) {
discard_fragment();
}
float4 srcColor = tint_symbol_7.sample(tint_symbol_8, texcoord);

4
test/bug/tint/749.spvasm.expected.msl
View File

@@ -1451,7 +1451,7 @@ void main_1(constant buf0& x_188, thread QuicksortObject* const tint_symbol_86,
int const x_909 = (*(tint_symbol_86)).numbers.arr[4];
(*(tint_symbol_86)).numbers.arr[4] = 0;
(*(tint_symbol_86)).numbers.arr[4] = x_909;
if (( fabs((x_308 - x_310)) < 0.25f)) {
if ((fabs((x_308 - x_310)) < 0.25f)) {
float const x_910 = uv.x;
uv.x = 0.0f;
uv.x = x_910;
@@ -1505,7 +1505,7 @@ void main_1(constant buf0& x_188, thread QuicksortObject* const tint_symbol_86,
uv[0] = 0.0f;
uv[0] = x_922;
float3 const x_519 = float3(x_447.x, x_446.x, x_446.y);
float3 const x_326 = normalize(x_325);
float3 const x_326 = normalize(x_325);
float const x_923 = uv.x;
uv.x = 0.0f;
uv.x = x_923;

4
test/bug/tint/913.wgsl.expected.msl
View File

@@ -13,7 +13,7 @@ struct OutputBuf {
};
bool aboutEqual(float value, float expect) {
return ( fabs((value - expect)) < 0.001f);
return (fabs((value - expect)) < 0.001f);
}
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture(0)]], texture2d<float, access::sample> tint_symbol_3 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
@@ -23,7 +23,7 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture
float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
if (((((dstTexCoord.x < uniforms.dstCopyOrigin.x) || (dstTexCoord.y < uniforms.dstCopyOrigin.y)) || (dstTexCoord.x >= (uniforms.dstCopyOrigin.x + uniforms.copySize.x))) || (dstTexCoord.y >= (uniforms.dstCopyOrigin.y + uniforms.copySize.y)))) {
success = (success && all((tint_symbol_3.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
success = (success && all((tint_symbol_3.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
} else {
uint2 srcTexCoord = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin);
if ((uniforms.dstTextureFlipY == 1u)) {

6
test/bug/tint/914.wgsl.expected.msl
View File

@@ -59,7 +59,7 @@ kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3
tint_array_wrapper const tint_symbol_3 = {.arr={}};
tint_symbol_5 = tint_symbol_3;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const tileRow = (local_id.y * RowPerThread);
uint const tileCol = (local_id.x * ColPerThread);
uint const globalRow = (global_id.y * RowPerThread);
@@ -140,7 +140,7 @@ kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3
}
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_threadgroup);
{
uint k = 0u;
while (true) {
@@ -189,7 +189,7 @@ kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3
}
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_threadgroup);
{
t = (t + 1u);
}