chore: push from SNAPKITTYWEST local build

ce25ab4 verified about 2 months ago

2.31 kB

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

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