File size: 2,115 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 | // SASS for dequant_gguf_q4k (sm_89)
// GGUF Q4_K block dequantization to FP16
// Author: Ahmad Ali Parr · Trust: Bel Esprit D'Accord Irrevocable Trust
dequant_gguf_q4k:
/* 0x0000 */ MOV R0, c[0x0][0x140]; // packed_ptr
/* 0x0008 */ MOV R1, c[0x0][0x148]; // scales_ptr
/* 0x0010 */ MOV R2, c[0x0][0x150]; // mins_ptr
/* 0x0018 */ MOV R3, c[0x0][0x158]; // output_ptr
/* 0x0020 */ MOV R4, c[0x0][0x160]; // num_blocks
/* 0x0028 */ S2R R5, SR_CTAID_X;
/* 0x0030 */ S2R R8, SR_TID_X;
/* 0x0038 */ IMAD R16, R5, 0x100, R8; // block_idx = blockIdx.x*256 + threadIdx.x
/* 0x0040 */ ISETP.GE.AND P0, PT, R16, R4, PT;
/* 0x0048 */ @P0 BRA EXIT;
/* 0x0050 */ IMAD R18, R16, 0x24, RZ; // byte offset to block
/* 0x0058 */ IADD3 R19, R0, R18, RZ; // packed data start
/* 0x0060 */ SHL.B32 R20, R16, 0x1; // scale/min offset (FP16)
/* 0x0068 */ IADD3 R20, R1, R20, RZ;
/* 0x0070 */ SHL.B32 R21, R16, 0x1;
/* 0x0078 */ IADD3 R21, R2, R21, RZ;
/* 0x0080 */ LDG.E.16 HR0, [R20]; // scale (FP16)
/* 0x0088 */ LDG.E.16 HR1, [R21]; // min (FP16)
/* 0x0090 */ CVT.F32.F16 FR0, HR0;
/* 0x0098 */ CVT.F32.F16 FR1, HR1;
/* 0x00a0 */ SHL.B32 R24, R8, 0x2; // thread byte offset (4 bytes)
/* 0x00a8 */ IADD3 R26, R19, R24, RZ;
/* 0x00b0 */ LDG.E.32 R27, [R26]; // 4 packed bytes
/* 0x00b8 */ MOV R28, 0x0; // value index
DEQUANT_LOOP:
/* 0x00c0 */ ISETP.GE.AND P1, PT, R28, 0x8, PT;
/* 0x00c8 */ @P1 BRA DEQUANT_DONE;
/* 0x00d0 */ SHL.B32 R29, R28, 0x2; // nibble shift
/* 0x00d8 */ SHR.B32 R30, R27, R29;
/* 0x00e0 */ AND.B32 R30, R30, 0xF; // extract nibble
/* 0x00e8 */ CVT.F32.U32 FR2, R30;
/* 0x00f0 */ FFMA FR2, FR2, FR0, FR1; // q*scale + min
/* 0x00f8 */ CVT.RNI.F16.F32 HR2, FR2;
/* 0x0100 */ IMAD R33, R16, 0x100, R8;// block base + thread
/* 0x0108 */ IMAD R33, R33, 0x8, R28; // + value index
/* 0x0110 */ SHL.B32 R33, R33, 0x1; // FP16 byte offset
/* 0x0118 */ IADD3 R37, R3, R33, RZ;
/* 0x0120 */ STG.E.16 [R37], HR2;
/* 0x0128 */ IADD3 R28, R28, 0x1, RZ;
/* 0x0130 */ BRA DEQUANT_LOOP;
DEQUANT_DONE:
EXIT:
/* 0x0138 */ EXIT;
|