mirror of
https://github.com/zenlm/engine.git
synced 2026-10-30 15:00:05 +00:00
Eric BuehlerandGitHub
1269bd8ab6
* Add the kernels * Remove include of aten or torch * Add the ffi bindings * Sketch the forward method * Handle input and output reshapes * Add some features * Improve compat of build.rs * Fix workspace dep * Finish merge * Fixes * Finish gptq gemm and add trait * Add the cuda gptq matmul stub * Remove default feature * Correct conditional comp * Add gguf qmatmul quantized support * Implement matmul with qmethod in qllama * Update readme of mistralrs quant * int* to int64* in q_gemm.cu * Add model and pipeline * Add gptq loader selector * Rename quantized_config -> quantization_config * Fix g_idx shape * Ensure WNA16 * Broadcast add * Format * int64_t* -> int* rollback * Prep for correct types * Finish merge * Complete merge * Update cargo lock * Integrate with new i32 type * Fixes * More progress * It doesnt crash * Oops * It works! * Remove some todos * Testing isq support * Add to all non adapter models * Clippy * Avoid reallocating * Add support for gptq to adapter models * Add docs and logging * Update docs * Clippy * Remove a todo
52 lines
1.4 KiB
Plaintext
52 lines
1.4 KiB
Plaintext
/*
|
|
Copied from https://github.com/turboderp/exllamav2
|
|
*/
|
|
|
|
#ifndef _qdq_util_cuh
|
|
#define _qdq_util_cuh
|
|
|
|
union half2_uint32 {
|
|
uint32_t as_uint32;
|
|
half2 as_half2;
|
|
__device__ half2_uint32(uint32_t val) : as_uint32(val) {}
|
|
__device__ half2_uint32(half2 val) : as_half2(val) {}
|
|
};
|
|
|
|
union half_uint16 {
|
|
uint16_t as_uint16;
|
|
half as_half;
|
|
__device__ half_uint16(uint16_t val) : as_uint16(val) {}
|
|
__device__ half_uint16(half val) : as_half(val) {}
|
|
};
|
|
|
|
// Max_scale premultiplied by 1/256
|
|
|
|
__forceinline__ __device__ half dq_scale(const int qs, const half max_scale) {
|
|
int qs_i = qs + 1;
|
|
half qs_h = __int2half_rn(qs_i * qs_i);
|
|
qs_h = __hmul(qs_h, max_scale);
|
|
return qs_h;
|
|
}
|
|
|
|
__forceinline__ __device__ half dq(const int q, const int qzero,
|
|
const half scale) {
|
|
return __hmul(__int2half_rn(q - qzero), scale);
|
|
}
|
|
|
|
__forceinline__ __device__ half dq_ns(const int q, const int qzero) {
|
|
// return __hsub(__int2half_rn(q), __int2half_rn(qzero));
|
|
return __int2half_rn(q - qzero);
|
|
}
|
|
|
|
__forceinline__ __device__ int exb(const uint32_t q, const int shift,
|
|
const int mask) {
|
|
return (int)((q >> shift) & mask);
|
|
}
|
|
|
|
__forceinline__ __device__ int exb(const uint32_t q1, const uint32_t q0,
|
|
const int shift, const int mask) {
|
|
return (int)(__funnelshift_rc(q0, q1, shift) & mask);
|
|
}
|
|
|
|
#endif
|