mirror of
https://github.com/zenlm/engine.git
synced 2026-10-30 15:00:05 +00:00
Eric BuehlerandGitHub
951871d010
* 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
25 lines
743 B
Plaintext
25 lines
743 B
Plaintext
/*
|
|
Copied from https://github.com/turboderp/exllamav2
|
|
*/
|
|
|
|
#ifndef _qdq_8_cuh
|
|
#define _qdq_8_cuh
|
|
|
|
#include "qdq_util.cuh"
|
|
|
|
__forceinline__ __device__ void shuffle_8bit_4(uint32_t* q, int stride) {}
|
|
|
|
__forceinline__ __device__ void dequant_8bit_8(const uint32_t q_0,
|
|
const uint32_t q_1,
|
|
half2 (&dq)[4], int stride,
|
|
const uint32_t zero) {
|
|
half dqh[8];
|
|
for (int i = 0; i < 4; i++) dqh[i] = dq_ns(exb(q_0, i * 8, 0xff), zero);
|
|
for (int i = 0; i < 4; i++) dqh[i + 4] = dq_ns(exb(q_1, i * 8, 0xff), zero);
|
|
|
|
for (int i = 0; i < 4; i++)
|
|
dq[i] = __halves2half2(dqh[i * 2], dqh[i * 2 + 1]);
|
|
}
|
|
|
|
#endif
|