File size: 2,314 Bytes
ce25ab4 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 | .version 8.5
.target sm_89
.address_size 64
// dequant.ptx — GGUF Block Quantization Dequantization (Q4_K)
// SnapKitty Sovereign Kernel — Hopper SM89
// Dequantizes 32-value blocks from GGUF format to FP16
// Author: Ahmad Ali Parr · Trust: Bel Esprit D'Accord Irrevocable Trust
.entry dequant_gguf_q4k(
.param .b64 packed_ptr,
.param .b64 scales_ptr,
.param .b64 mins_ptr,
.param .b64 output_ptr,
.param .b32 num_blocks
)
{
.reg .pred %p<32>;
.reg .b32 %r<64>;
.reg .b64 %rd<32>;
.reg .f16 %h<128>;
.reg .f32 %f<64>;
ld.param.u64 %rd0, [packed_ptr];
ld.param.u64 %rd1, [scales_ptr];
ld.param.u64 %rd2, [mins_ptr];
ld.param.u64 %rd3, [output_ptr];
ld.param.u32 %r0, [num_blocks];
mov.u32 %r1, %ctaid.x;
mov.u32 %r4, %tid.x;
// Global block index
mul.lo.u32 %r10, %r1, 256;
add.u32 %r10, %r10, %r4;
// Bounds check
setp.ge.u32 %p0, %r10, %r0;
@%p0 bra EXIT;
// Block offset: 36 bytes per block
mul.lo.u32 %r11, %r10, 36;
cvt.u64.u32 %rd4, %r11;
add.u64 %rd4, %rd0, %rd4;
// Scale and min offset
cvt.u64.u32 %rd5, %r10;
mul.lo.u64 %rd5, %rd5, 2; // FP16 = 2 bytes
add.u64 %rd5, %rd1, %rd5;
cvt.u64.u32 %rd6, %r10;
mul.lo.u64 %rd6, %rd6, 2;
add.u64 %rd6, %rd2, %rd6;
ld.global.ca.f16 %h0, [%rd5];
ld.global.ca.f16 %h1, [%rd6];
cvt.f32.f16 %f0, %h0; // scale
cvt.f32.f16 %f1, %h1; // min
// Thread byte offset (8 values per thread = 4 bytes)
mul.lo.u32 %r12, %r4, 4;
cvt.u64.u32 %rd7, %r12;
add.u64 %rd7, %rd4, %rd7;
ld.global.ca.b32 %r15, [%rd7];
mov.u32 %r16, 0;
DEQUANT_LOOP:
setp.ge.u32 %p1, %r16, 8;
@%p1 bra DEQUANT_DONE;
// Extract nibble
shl.b32 %r17, %r16, 2;
shr.b32 %r18, %r15, %r17;
and.b32 %r18, %r18, 0xF;
// Dequantize: q * scale + min
cvt.rn.f32.u32 %f2, %r18;
fma.rn.f32 %f3, %f2, %f0, %f1;
cvt.rn.f16.f32 %h2, %f3;
// Store
mul.lo.u32 %r21, %r10, 256;
add.u32 %r21, %r21, %r4;
mul.lo.u32 %r21, %r21, 8;
add.u32 %r21, %r21, %r16;
mul.lo.u32 %r21, %r21, 2;
cvt.u64.u32 %rd8, %r21;
add.u64 %rd8, %rd3, %rd8;
st.global.ca.f16 [%rd8], %h2;
add.u32 %r16, %r16, 1;
bra DEQUANT_LOOP;
DEQUANT_DONE:
EXIT:
ret;
}
|