| 22 | const std::string versionToStr(int version); |
| 23 | |
| 24 | void matmulf16_forward_cpu(half* out, |
| 25 | const half* inp, const half* weight, const half* bias, |
| 26 | int B, int T, int C, int OC) { |
| 27 | // OC is short for "output channels" |
| 28 | // inp is (B,T,C), weight is (OC, C) |
| 29 | // out will be (B,T,OC) |
| 30 | #pragma omp parallel for collapse(2) |
| 31 | for (int b = 0; b < B; b++) { |
| 32 | for (int t = 0; t < T; t++) { |
| 33 | half* out_bt = out + b * T * OC + t * OC; |
| 34 | const half* inp_bt = inp + b * T * C + t * C; |
| 35 | for (int o = 0; o < OC; o++) { |
| 36 | float val = (bias != NULL) ? halfToFloat(bias[o]) : 0.0f; |
| 37 | const half* wrow = weight + o*C; |
| 38 | for (int i = 0; i < C; i++) { |
| 39 | val += halfToFloat(inp_bt[i]) * halfToFloat(wrow[i]); |
| 40 | } |
| 41 | out_bt[o] = val; |
| 42 | } |
| 43 | } |
| 44 | } |
| 45 | } |
| 46 | |
| 47 | static const char *kShaderMatmul1 = R"( |
| 48 | @group(0) @binding(0) var<storage, read_write> A: array<{{precision}}>; |
no test coverage detected