mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-08-06 12:15:43 +00:00
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>
50 lines
1.4 KiB
Plaintext
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;
|
|
}
|
|
|