.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;
}