mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-08-05 03:35:56 +00:00
663271dca4
Inline the `continuing` block in the places where `continue` is called. Simplifies the emission, and fixes emission of `let` statements in the loop. This fix matches the same approach in writer/hlsl. See: https://dawn-review.googlesource.com/c/tint/+/51784 Fixed: tint:833 Fixed: tint:914 Change-Id: If4d8cde62dfaf8efa24272854ca7ff5edc0a8234 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55341 Commit-Queue: Ben Clayton <bclayton@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: David Neto <dneto@google.com>
37 lines
1.1 KiB
Plaintext
37 lines
1.1 KiB
Plaintext
#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];
|
|
};
|
|
|
|
kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
|
|
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;
|
|
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);
|
|
}
|
|
}
|
|
}
|
|
uint const index = (resultCell.y + (resultCell.x * dimOutter));
|
|
resultMatrix.numbers[index] = result;
|
|
return;
|
|
}
|
|
|