xref: /aosp_15_r20/external/skia/tests/sksl/compute/MatrixMultiply.wgsl (revision c8dee2aa9b3f27cf6c858bd81872bdeb2c07ed17)
1*c8dee2aaSAndroid Build Coastguard Workerdiagnostic(off, derivative_uniformity);
2*c8dee2aaSAndroid Build Coastguard Workerdiagnostic(off, chromium.unreachable_code);
3*c8dee2aaSAndroid Build Coastguard Workerstruct CSIn {
4*c8dee2aaSAndroid Build Coastguard Worker  @builtin(global_invocation_id) sk_GlobalInvocationID: vec3<u32>,
5*c8dee2aaSAndroid Build Coastguard Worker};
6*c8dee2aaSAndroid Build Coastguard Workerstruct sizeBuffer {
7*c8dee2aaSAndroid Build Coastguard Worker  sizes: array<vec2<i32>>,
8*c8dee2aaSAndroid Build Coastguard Worker};
9*c8dee2aaSAndroid Build Coastguard Worker@group(0) @binding(0) var<storage, read_write> _storage0 : sizeBuffer;
10*c8dee2aaSAndroid Build Coastguard Workerstruct inputs1 {
11*c8dee2aaSAndroid Build Coastguard Worker  data1: array<f32>,
12*c8dee2aaSAndroid Build Coastguard Worker};
13*c8dee2aaSAndroid Build Coastguard Worker@group(0) @binding(1) var<storage, read> _storage1 : inputs1;
14*c8dee2aaSAndroid Build Coastguard Workerstruct inputs2 {
15*c8dee2aaSAndroid Build Coastguard Worker  data2: array<f32>,
16*c8dee2aaSAndroid Build Coastguard Worker};
17*c8dee2aaSAndroid Build Coastguard Worker@group(0) @binding(2) var<storage, read> _storage2 : inputs2;
18*c8dee2aaSAndroid Build Coastguard Workerstruct result {
19*c8dee2aaSAndroid Build Coastguard Worker  resultData: array<f32>,
20*c8dee2aaSAndroid Build Coastguard Worker};
21*c8dee2aaSAndroid Build Coastguard Worker@group(0) @binding(3) var<storage, read_write> _storage3 : result;
22*c8dee2aaSAndroid Build Coastguard Workerfn _skslMain(_stageIn: CSIn) {
23*c8dee2aaSAndroid Build Coastguard Worker  {
24*c8dee2aaSAndroid Build Coastguard Worker    _storage0.sizes[2] = vec2<i32>(_storage0.sizes[0].x, _storage0.sizes[1].y);
25*c8dee2aaSAndroid Build Coastguard Worker    let resultCell: vec2<i32> = vec2<i32>(i32(_stageIn.sk_GlobalInvocationID.x), i32(_stageIn.sk_GlobalInvocationID.y));
26*c8dee2aaSAndroid Build Coastguard Worker    var result: f32 = 0.0;
27*c8dee2aaSAndroid Build Coastguard Worker    {
28*c8dee2aaSAndroid Build Coastguard Worker      var i: i32 = 0;
29*c8dee2aaSAndroid Build Coastguard Worker      loop {
30*c8dee2aaSAndroid Build Coastguard Worker        if i < _storage0.sizes[0].y {
31*c8dee2aaSAndroid Build Coastguard Worker          {
32*c8dee2aaSAndroid Build Coastguard Worker            let a: i32 = i + resultCell.x * _storage0.sizes[0].y;
33*c8dee2aaSAndroid Build Coastguard Worker            let b: i32 = resultCell.y + i * _storage0.sizes[1].y;
34*c8dee2aaSAndroid Build Coastguard Worker            result = result + _storage1.data1[a] * _storage2.data2[b];
35*c8dee2aaSAndroid Build Coastguard Worker          }
36*c8dee2aaSAndroid Build Coastguard Worker        } else {
37*c8dee2aaSAndroid Build Coastguard Worker          break;
38*c8dee2aaSAndroid Build Coastguard Worker        }
39*c8dee2aaSAndroid Build Coastguard Worker        continuing {
40*c8dee2aaSAndroid Build Coastguard Worker          i = i + i32(1);
41*c8dee2aaSAndroid Build Coastguard Worker        }
42*c8dee2aaSAndroid Build Coastguard Worker      }
43*c8dee2aaSAndroid Build Coastguard Worker    }
44*c8dee2aaSAndroid Build Coastguard Worker    let index: i32 = resultCell.y + resultCell.x * _storage0.sizes[1].y;
45*c8dee2aaSAndroid Build Coastguard Worker    _storage3.resultData[index] = result;
46*c8dee2aaSAndroid Build Coastguard Worker  }
47*c8dee2aaSAndroid Build Coastguard Worker}
48*c8dee2aaSAndroid Build Coastguard Worker@compute @workgroup_size(16, 16, 1) fn main(_stageIn: CSIn) {
49*c8dee2aaSAndroid Build Coastguard Worker  _skslMain(_stageIn);
50*c8dee2aaSAndroid Build Coastguard Worker}
51