mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-08-06 12:15:43 +00:00
7103f51603
Metal 1.x does not support swizzling on packed_vec types. Use array-index for single element selection (permitted on LHS and RHS of assignment) Cast the packed_vec to a vec for multiple element swizzles (not permitted as the LHS of an assignment). Fixed: tint:1249 Change-Id: I70cbb0c22a935b06b3905d24484bdc2edfb95fc2 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67060 Commit-Queue: Ben Clayton <bclayton@google.com> Kokoro: Ben Clayton <bclayton@google.com> Reviewed-by: James Price <jrprice@google.com>
32 lines
1.2 KiB
Plaintext
32 lines
1.2 KiB
Plaintext
#include <metal_stdlib>
|
|
|
|
using namespace metal;
|
|
struct Uniforms {
|
|
/* 0x0000 */ uint2 aShape;
|
|
/* 0x0008 */ uint2 bShape;
|
|
/* 0x0010 */ uint2 outShape;
|
|
};
|
|
struct Matrix {
|
|
/* 0x0000 */ uint numbers[1];
|
|
};
|
|
|
|
void tint_symbol_inner(constant Uniforms& uniforms, const device Matrix& firstMatrix, const device Matrix& secondMatrix, device Matrix& resultMatrix, uint3 global_id) {
|
|
uint2 const resultCell = uint2(global_id[1], global_id[0]);
|
|
uint const dimInner = uniforms.aShape[1];
|
|
uint const dimOutter = uniforms.outShape[1];
|
|
uint result = 0u;
|
|
for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
|
|
uint const a = (i + (resultCell[0] * dimInner));
|
|
uint const b = (resultCell[1] + (i * dimOutter));
|
|
result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
|
|
}
|
|
uint const index = (resultCell[1] + (resultCell[0] * dimOutter));
|
|
resultMatrix.numbers[index] = result;
|
|
}
|
|
|
|
kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(0)]], const device Matrix& firstMatrix [[buffer(2)]], const device Matrix& secondMatrix [[buffer(3)]], device Matrix& resultMatrix [[buffer(1)]]) {
|
|
tint_symbol_inner(uniforms, firstMatrix, secondMatrix, resultMatrix, global_id);
|
|
return;
|
|
}
|
|
|