| 1 | /* |
| 2 | Copied from https://github.com/turboderp/exllamav2 |
| 3 | */ |
| 4 | |
| 5 | #ifndef _qdq_8_cuh |
| 6 | #define _qdq_8_cuh |
| 7 | |
| 8 | #include "qdq_util.cuh" |
| 9 | |
| 10 | namespace vllm { |
| 11 | namespace gptq { |
| 12 | |
| 13 | __forceinline__ __device__ void shuffle_8bit_4(uint32_t* q, int stride) {} |
| 14 | |
| 15 | __forceinline__ __device__ void dequant_8bit_8(const uint32_t q_0, |
| 16 | const uint32_t q_1, |
| 17 | half2 (&dq)[4], int stride, |
| 18 | const uint32_t zero) { |
| 19 | half dqh[8]; |
| 20 | for (int i = 0; i < 4; i++) dqh[i] = dq_ns(exb(q_0, i * 8, 0xff), zero); |
| 21 | for (int i = 0; i < 4; i++) dqh[i + 4] = dq_ns(exb(q_1, i * 8, 0xff), zero); |
| 22 | |
| 23 | for (int i = 0; i < 4; i++) |
| 24 | dq[i] = __halves2half2(dqh[i * 2], dqh[i * 2 + 1]); |
| 25 | } |
| 26 | |
| 27 | } // namespace gptq |
| 28 | } // namespace vllm |
| 29 | |
| 30 | #endif |