| 32 | // queue: the SYCL queue to submit into |
| 33 | // ────────────────────────────────────────────────────────────────────────────── |
| 34 | extern "C" ESIMD_KERNEL_API void sdp_fp16( |
| 35 | void* Q, void* K, void* V, |
| 36 | void* normAlpha, |
| 37 | void* out, |
| 38 | int q_len, int kv_len, |
| 39 | int headQ, int headKv, |
| 40 | void* sycl_queue_ptr) |
| 41 | { |
| 42 | sycl::queue& q = *reinterpret_cast<sycl::queue*>(sycl_queue_ptr); |
| 43 | |
| 44 | int groupH = headQ; |
| 45 | int groupV = (q_len + 255) / 256; |
| 46 | sycl::nd_range<2> ndr({(size_t)(16 * groupH), (size_t)groupV}, {16, 1}); |
| 47 | |
| 48 | uint8_t* pQ = reinterpret_cast<uint8_t*>(Q); |
| 49 | uint8_t* pK = reinterpret_cast<uint8_t*>(K); |
| 50 | uint8_t* pV = reinterpret_cast<uint8_t*>(V); |
| 51 | uint8_t* pA = reinterpret_cast<uint8_t*>(normAlpha); |
| 52 | uint8_t* pO = reinterpret_cast<uint8_t*>(out); |
| 53 | uint32_t aLen = (uint32_t)q_len; |
| 54 | uint32_t kvLen = (uint32_t)kv_len; |
| 55 | uint32_t hQ = (uint32_t)headQ; |
| 56 | uint32_t hKv = (uint32_t)headKv; |
| 57 | |
| 58 | q.submit([&](sycl::handler& cgh) { |
| 59 | cgh.parallel_for(ndr, [=](sycl::nd_item<2> ndi) SYCL_ESIMD_KERNEL { |
| 60 | flashAttnBMha128Fp16OptPrecomputed( |
| 61 | pQ, pK, pV, pA, pO, |
| 62 | aLen, kvLen, hQ, hKv, ndi); |
| 63 | }); |
| 64 | }).wait(); |
| 65 | } |
| 66 | |
| 67 | // ────────────────────────────────────────────────────────────────────────────── |
| 68 | // sdp_bf16io: BF16 I/O hybrid Flash Attention |
nothing calls this directly
no test coverage detected