dawn-cmake/test/bug/tint/744.wgsl.expected.msl
James Price 7697c31e84 writer/msl: Emit builtins as parameters
Add a config parameter for the CanonicalizeEntryPoint transform that
selects between emitting builtins as parameters (for MSL) or struct
members (for HLSL).

This fixes all of the shader IO issues in Tint's E2E tests for MSL.

Fixed: tint:817
Change-Id: Ieb31cdbd2e4d96ac41f8d8515fd07ead8241d770
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53282
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
2021-06-04 14:40:28 +00:00

50 lines
1.4 KiB
Plaintext

SKIP: crbug.com/tint/833 loop emission is broken
#include <metal_stdlib>
using namespace metal;
struct Uniforms {
/* 0x0000 */ packed_uint2 aShape;
/* 0x0008 */ packed_uint2 bShape;
/* 0x0010 */ packed_uint2 outShape;
};
struct Matrix {
/* 0x0000 */ uint numbers[1];
};
struct tint_symbol_2 {
uint3 global_id [[thread_position_in_grid]];
};
kernel void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
uint3 const global_id = tint_symbol_1.global_id;
uint2 const resultCell = uint2(global_id.y, global_id.x);
uint const dimInner = uniforms.aShape.y;
uint const dimOutter = uniforms.outShape.y;
uint result = 0u;
{
uint i = 0u;
{
bool tint_msl_is_first_1 = true;
uint const a;
uint const b;
for(;;) {
if (!tint_msl_is_first_1) {
i = (i + 1u);
}
tint_msl_is_first_1 = false;
if (!((i < dimInner))) {
break;
}
a = (i + (resultCell.x * dimInner));
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;
return;
}