| 1 | /* |
| 2 | Copied from https://github.com/turboderp/exllamav2 |
| 3 | */ |
| 4 | |
| 5 | #ifndef _qdq_util_cuh |
| 6 | #define _qdq_util_cuh |
| 7 | |
| 8 | namespace vllm { |
| 9 | namespace gptq { |
| 10 | |
| 11 | union half2_uint32 { |
| 12 | uint32_t as_uint32; |
| 13 | half2 as_half2; |
| 14 | __device__ half2_uint32(uint32_t val) : as_uint32(val) {} |
| 15 | __device__ half2_uint32(half2 val) : as_half2(val) {} |
| 16 | }; |
| 17 | |
| 18 | union half_uint16 { |
| 19 | uint16_t as_uint16; |
| 20 | half as_half; |
| 21 | __device__ half_uint16(uint16_t val) : as_uint16(val) {} |
| 22 | __device__ half_uint16(half val) : as_half(val) {} |
| 23 | }; |
| 24 | |
| 25 | __forceinline__ __device__ half dq(const int q, const int qzero, |
| 26 | const half scale) { |
| 27 | return __hmul(__int2half_rn(q - qzero), scale); |
| 28 | } |
| 29 | |
| 30 | __forceinline__ __device__ half dq_ns(const int q, const int qzero) { |
| 31 | // return __hsub(__int2half_rn(q), __int2half_rn(qzero)); |
| 32 | return __int2half_rn(q - qzero); |
| 33 | } |
| 34 | |
| 35 | __forceinline__ __device__ int exb(const uint32_t q, const int shift, |
| 36 | const int mask) { |
| 37 | return (int)((q >> shift) & mask); |
| 38 | } |
| 39 | |
| 40 | __forceinline__ __device__ int exb(const uint32_t q1, const uint32_t q0, |
| 41 | const int shift, const int mask) { |
| 42 | return (int)(__funnelshift_rc(q0, q1, shift) & mask); |
| 43 | } |
| 44 | |
| 45 | } // namespace gptq |
| 46 | } // namespace vllm |
| 47 | #endif |