Skip to content
KernelIndex
Search⌘K
Sign in

submission 728835

bhagawan-yantrion · python · License unknown
Kernel source · 707 lines ↓holds 1 record

Use it

Vendorable · source mirrored · license unknownView source →

No package. Vendor the mirrored source: 707 lines, June 9 Researcher Reciprocity License v1.0.

submission_yantrion.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-amd-mxfp4-mm-728835?include=source"
interfacepython
Compatibility
measured onAMD Instinct MI355X
declared hardwareAMD Instinct MI355X
architecturesgfx950
dtypesbf16, mxfp4

Benchmark evidence

1 measurement across 1 GPU, fastest first.

Operation / workload
Hardware
Latency
Rank
Observed
AMD MXFP4 GEMMsuite of 6 cases
AMD Instinct MI355X
4.35µs
#1 of 1143
2026-04-04
Run →

Reported · How evidence levels are derived →

Source and license

sourceavailable
revision digestsha256:666a1da9e01facc51656c7b9a1064365f9edd19e05748bd04ce2891009d6df07
license declaredunknown
license concludedunknown
authorsbhagawan-yantrion
imported2026-08-15

Techniques

Extracted from the mirrored source by pattern, never inferred. Each row cites its line.

split-kdef _get_split_k(m, n, k):

Kernel source

submission_yantrion.py707 lines
#!/usr/bin/env python3
"""HSA AQL direct dispatch — single-file submission for testing.
Falls back to hipModuleLaunchKernel if AQL init fails."""
import os, ctypes, base64, tempfile, struct, sys
os.environ["HIP_FORCE_DEV_KERNARG"] = "1"
import torch
from task import input_t, output_t

# ═══ HSA Constants ═══
_HSA_OK = 0
_HSA_BREAK = 1
_HSA_AGENT_DEVICE = 17
_HSA_AGENT_QMAX = 14
_HSA_DEV_GPU = 1
_HSA_Q_SINGLE = 0
_HSA_PKT_DISPATCH = 2
_HSA_PKT_INVALID = 0
_HSA_PKT_BARRIER_AND = 3
_HSA_FENCE_AGENT = 2
_HSA_COND_LT = 1
_HSA_WAIT_ACTIVE = 1
_HSA_PROFILE_FULL = 1
_HSA_ROUND_DEFAULT = 0
_HSA_SYM_KERNEL_OBJ = 22
_HSA_SYM_PRIV_SEG = 25
_HSA_SYM_GROUP_SEG = 26
_HSA_REG_SEGMENT = 0
_HSA_REG_GFLAGS = 1
_HSA_SEG_GLOBAL = 0
_HSA_FLAG_KERNARG = 8

# ═══ HSA ctypes structs ═══
class _ag(ctypes.Structure):
    _fields_ = [("handle", ctypes.c_uint64)]
class _sig(ctypes.Structure):
    _fields_ = [("handle", ctypes.c_uint64)]
class _q(ctypes.Structure):
    _fields_ = [("type", ctypes.c_uint32), ("features", ctypes.c_uint32),
                ("base_address", ctypes.c_void_p), ("doorbell_signal", _sig),
                ("size", ctypes.c_uint32), ("reserved1", ctypes.c_uint32),
                ("id", ctypes.c_uint64)]
class _pkt(ctypes.Structure):
    _fields_ = [("header", ctypes.c_uint16), ("setup", ctypes.c_uint16),
                ("wg_x", ctypes.c_uint16), ("wg_y", ctypes.c_uint16),
                ("wg_z", ctypes.c_uint16), ("reserved0", ctypes.c_uint16),
                ("grid_x", ctypes.c_uint32), ("grid_y", ctypes.c_uint32),
                ("grid_z", ctypes.c_uint32),
                ("private_seg", ctypes.c_uint32), ("group_seg", ctypes.c_uint32),
                ("kernel_object", ctypes.c_uint64), ("kernarg_address", ctypes.c_uint64),
                ("reserved2", ctypes.c_uint64), ("completion_signal", _sig)]

_AGENT_CB = ctypes.CFUNCTYPE(ctypes.c_uint32, _ag, ctypes.c_void_p)
_REGION_CB = ctypes.CFUNCTYPE(ctypes.c_uint32, ctypes.c_uint64, ctypes.c_void_p)

def _mkhdr(pt, bar=1, acq=_HSA_FENCE_AGENT, rel=_HSA_FENCE_AGENT):
    return pt | (bar << 8) | (acq << 9) | (rel << 11)

# ═══ Kernel binary ═══
_CO_B64 = (
    "f0VMRgIBAUAEAAAAAAAAAAMA4AABAAAAAAAAAAAAAABAAAAAAAAAANgrAAAAAAAATwUAAEAAOAAJAEAAEAAOAAYAAAAEAAAAQAAA"
    "AAAAAABAAAAAAAAAAEAAAAAAAAAA+AEAAAAAAAD4AQAAAAAAAAgAAAAAAAAAAQAAAAQAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AABADgAAAAAAAEAOAAAAAAAAABAAAAAAAAABAAAABQAAAAAPAAAAAAAAAB8AAAAAAAAAHwAAAAAAAMAVAAAAAAAAwBUAAAAAAAAA"
    "EAAAAAAAAAEAAAAGAAAAwCQAAAAAAADARAAAAAAAAMBEAAAAAAAAcAAAAAAAAABACwAAAAAAAAAQAAAAAAAAAQAAAAYAAAAwJQAA"
    "AAAAADBVAAAAAAAAMFUAAAAAAAAAAAAAAAAAAAEAAAAAAAAAABAAAAAAAAACAAAABgAAAMAkAAAAAAAAwEQAAAAAAADARAAAAAAA"
    "AHAAAAAAAAAAcAAAAAAAAAAIAAAAAAAAAFLldGQEAAAAwCQAAAAAAADARAAAAAAAAMBEAAAAAAAAcAAAAAAAAABACwAAAAAAAAEA"
    "AAAAAAAAUeV0ZAYAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAEAAAABAAAADgCAAAA"
    "AAAAOAIAAAAAAAA4AgAAAAAAAPgJAAAAAAAA+AkAAAAAAAAEAAAAAAAAAAcAAADjCQAAIAAAAEFNREdQVQAAg65hbWRoc2Eua2Vy"
    "bmVsc5LeABKrLmFncHJfY291bnQApS5hcmdz3AAbha4uYWN0dWFsX2FjY2Vzc6lyZWFkX29ubHmuLmFkZHJlc3Nfc3BhY2WmZ2xv"
    "YmFspy5vZmZzZXQApS5zaXplCKsudmFsdWVfa2luZK1nbG9iYWxfYnVmZmVyha4uYWN0dWFsX2FjY2Vzc6lyZWFkX29ubHmuLmFk"
    "ZHJlc3Nfc3BhY2WmZ2xvYmFspy5vZmZzZXQIpS5zaXplCKsudmFsdWVfa2luZK1nbG9iYWxfYnVmZmVyha4uYWN0dWFsX2FjY2Vz"
    "c6lyZWFkX29ubHmuLmFkZHJlc3Nfc3BhY2WmZ2xvYmFspy5vZmZzZXQQpS5zaXplCKsudmFsdWVfa2luZK1nbG9iYWxfYnVmZmVy"
    "hK4uYWRkcmVzc19zcGFjZaZnbG9iYWynLm9mZnNldBilLnNpemUIqy52YWx1ZV9raW5krWdsb2JhbF9idWZmZXKDpy5vZmZzZXQg"
    "pS5zaXplBKsudmFsdWVfa2luZKhieV92YWx1ZYOnLm9mZnNldCSlLnNpemUEqy52YWx1ZV9raW5kqGJ5X3ZhbHVlg6cub2Zmc2V0"
    "KKUuc2l6ZQSrLnZhbHVlX2tpbmSoYnlfdmFsdWWDpy5vZmZzZXQspS5zaXplBKsudmFsdWVfa2luZKhieV92YWx1ZYOnLm9mZnNl"
    "dDClLnNpemUEqy52YWx1ZV9raW5kqGJ5X3ZhbHVlg6cub2Zmc2V0NKUuc2l6ZQSrLnZhbHVlX2tpbmSoYnlfdmFsdWWEri5hZGRy"
    "ZXNzX3NwYWNlpmdsb2JhbKcub2Zmc2V0OKUuc2l6ZQirLnZhbHVlX2tpbmStZ2xvYmFsX2J1ZmZlcoSuLmFkZHJlc3Nfc3BhY2Wm"
    "Z2xvYmFspy5vZmZzZXRApS5zaXplCKsudmFsdWVfa2luZK1nbG9iYWxfYnVmZmVyg6cub2Zmc2V0SKUuc2l6ZQSrLnZhbHVlX2tp"
    "bmSoYnlfdmFsdWWDpy5vZmZzZXRQpS5zaXplBKsudmFsdWVfa2luZLRoaWRkZW5fYmxvY2tfY291bnRfeIOnLm9mZnNldFSlLnNp"
    "emUEqy52YWx1ZV9raW5ktGhpZGRlbl9ibG9ja19jb3VudF95g6cub2Zmc2V0WKUuc2l6ZQSrLnZhbHVlX2tpbmS0aGlkZGVuX2Js"
    "b2NrX2NvdW50X3qDpy5vZmZzZXRcpS5zaXplAqsudmFsdWVfa2luZLNoaWRkZW5fZ3JvdXBfc2l6ZV94g6cub2Zmc2V0XqUuc2l6"
    "ZQKrLnZhbHVlX2tpbmSzaGlkZGVuX2dyb3VwX3NpemVfeYOnLm9mZnNldGClLnNpemUCqy52YWx1ZV9raW5ks2hpZGRlbl9ncm91"
    "cF9zaXplX3qDpy5vZmZzZXRipS5zaXplAqsudmFsdWVfa2luZLJoaWRkZW5fcmVtYWluZGVyX3iDpy5vZmZzZXRkpS5zaXplAqsu"
    "dmFsdWVfa2luZLJoaWRkZW5fcmVtYWluZGVyX3mDpy5vZmZzZXRmpS5zaXplAqsudmFsdWVfa2luZLJoaWRkZW5fcmVtYWluZGVy"
    "X3qDpy5vZmZzZXR4pS5zaXplCKsudmFsdWVfa2luZLZoaWRkZW5fZ2xvYmFsX29mZnNldF94g6cub2Zmc2V0zIClLnNpemUIqy52"
    "YWx1ZV9raW5ktmhpZGRlbl9nbG9iYWxfb2Zmc2V0X3mDpy5vZmZzZXTMiKUuc2l6ZQirLnZhbHVlX2tpbmS2aGlkZGVuX2dsb2Jh"
    "bF9vZmZzZXRfeoOnLm9mZnNldMyQpS5zaXplAqsudmFsdWVfa2luZLBoaWRkZW5fZ3JpZF9kaW1zg6cub2Zmc2V0zMilLnNpemUE"
    "qy52YWx1ZV9raW5kt2hpZGRlbl9keW5hbWljX2xkc19zaXpluS5ncm91cF9zZWdtZW50X2ZpeGVkX3NpemUEti5rZXJuYXJnX3Nl"
    "Z21lbnRfYWxpZ24ItS5rZXJuYXJnX3NlZ21lbnRfc2l6Zc0BUKkubGFuZ3VhZ2WoT3BlbkNMIEOxLmxhbmd1YWdlX3ZlcnNpb26S"
    "AgC4Lm1heF9mbGF0X3dvcmtncm91cF9zaXplzQIApS5uYW1lt2Z1c2VkX214ZnA0X2dlbW1fa2VybmVsuy5wcml2YXRlX3NlZ21l"
    "bnRfZml4ZWRfc2l6ZQCrLnNncHJfY291bnQ0sS5zZ3ByX3NwaWxsX2NvdW50AKcuc3ltYm9sumZ1c2VkX214ZnA0X2dlbW1fa2Vy"
    "bmVsLmtkuC51bmlmb3JtX3dvcmtfZ3JvdXBfc2l6ZQGzLnVzZXNfZHluYW1pY19zdGFja8KrLnZncHJfY291bnQ3sS52Z3ByX3Nw"
    "aWxsX2NvdW50AK8ud2F2ZWZyb250X3NpemVA3gASqy5hZ3ByX2NvdW50AKUuYXJnc5SFri5hY3R1YWxfYWNjZXNzqXJlYWRfb25s"
    "ea4uYWRkcmVzc19zcGFjZaZnbG9iYWynLm9mZnNldAClLnNpemUIqy52YWx1ZV9raW5krWdsb2JhbF9idWZmZXKFri5hY3R1YWxf"
    "YWNjZXNzqndyaXRlX29ubHmuLmFkZHJlc3Nfc3BhY2WmZ2xvYmFspy5vZmZzZXQIpS5zaXplCKsudmFsdWVfa2luZK1nbG9iYWxf"
    "YnVmZmVyg6cub2Zmc2V0EKUuc2l6ZQSrLnZhbHVlX2tpbmSoYnlfdmFsdWWDpy5vZmZzZXQUpS5zaXplBKsudmFsdWVfa2luZKhi"
    "eV92YWx1ZbkuZ3JvdXBfc2VnbWVudF9maXhlZF9zaXplALYua2VybmFyZ19zZWdtZW50X2FsaWduCLUua2VybmFyZ19zZWdtZW50"
    "X3NpemUYqS5sYW5ndWFnZahPcGVuQ0wgQ7EubGFuZ3VhZ2VfdmVyc2lvbpICALgubWF4X2ZsYXRfd29ya2dyb3VwX3NpemXNAQCl"
    "Lm5hbWWycmVkdWNlX2JmMTZfa2VybmVsuy5wcml2YXRlX3NlZ21lbnRfZml4ZWRfc2l6ZQCrLnNncHJfY291bnQMsS5zZ3ByX3Nw"
    "aWxsX2NvdW50AKcuc3ltYm9stXJlZHVjZV9iZjE2X2tlcm5lbC5rZLgudW5pZm9ybV93b3JrX2dyb3VwX3NpemUBsy51c2VzX2R5"
    "bmFtaWNfc3RhY2vCqy52Z3ByX2NvdW50BrEudmdwcl9zcGlsbF9jb3VudACvLndhdmVmcm9udF9zaXplQK1hbWRoc2EudGFyZ2V0"
    "uWFtZGdjbi1hbWQtYW1kaHNhLS1nZng5NTCuYW1kaHNhLnZlcnNpb26SAQIAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAQAAABID"
    "BwAAHwAAAAAAAEgQAAAAAAAAGQAAABEDBgDADQAAAAAAAEAAAAAAAAAANAAAABIDBwAAMAAAAAAAAKgAAAAAAAAARwAAABEDBgAA"
    "DgAAAAAAAEAAAAAAAAAAXQAAABEACgAwVQAAAAAAAAEAAAAAAAAAAQAAAAEAAAABAAAAGgAAAAgAFkCAmQAAAQAAAO6ca0ps/JkM"
    "qqvFeqhhyJ7R76FRBgAAAAYAAAAEAAAABQAAAAAAAAAAAAAAAgAAAAAAAAAAAAAAAAAAAAEAAAAAAAAAAwAAAAAAAAAAZnVzZWRf"
    "bXhmcDRfZ2VtbV9rZXJuZWwAZnVzZWRfbXhmcDRfZ2VtbV9rZXJuZWwua2QAcmVkdWNlX2JmMTZfa2VybmVsAHJlZHVjZV9iZjE2"
    "X2tlcm5lbC5rZABfX2hpcF9jdWlkXzk5OTU2YjkxZTQ4MjI1ZTkAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAQAAAAA"
    "AAAAUAEAAAAAAABAEQAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAADQAAAIYBrwCEAgAACAAAAAAAAAAAAAAAAAAAABgAAAAAAAAA"
    "ACIAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAEAAABAAK8AhAAAAAgAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AAAAAAAAAAAAAAAAAAAAAAEOwCAAAACAAwLAUAAAAIAHAsBYAAAAAp8PkA+dD49/wIy/Bf8MgX8AAAAMnw2QDZkNjwwNDIEMhyKQ"
    "Dp8MkAydDI8ODAyBAg8PgQyDDZAMyAyGD4MQkA/ID4YODIyBAg+PgQ8MDIMNDw2SDBAMgQwNDIEMDgS/DAIWhSIwgr4CDAJ+gAKO"
    "gRYiDYgNnx+QAUcCfhYwjb4ABgbAQAAAAAADAsBIAAAAAAUGwBgAAAD/AgIK/v9/TwEPAn5/wIy/DAMXkgEFHn4ODw6SDw4Olg8O"
    "D4ENDw6WDgIPkg0PjYEOgQ+BDQKQgQ0CCb8PDg6FEA0NhQ6BD4ENAgm/Dw4ChRcMDIEMBgyDDBecgRyfDI8cDAyBDIEkkByfDJAM"
    "mwyPHAwMgQIfIYgMhSWQIR+mgSSEHY4lhCeOJoQCjh2EI4EnAIh9aiCMvkkBiL8lMKi+KAwCfoADBsAAAAAAgCiTgSWFEo4BRwJ+"
    "JZ8pkCQnqoESAP3RAAkRAv8CAgr+/39PAQ8CfhQA/dEAC10AgAGQvoASq4ECAIXSEwICAAIAhtIBBQIAAQUCaP8ArL4ABQwM/wCt"
    "vgAA/wCAAiZ+lwIqfgADLH5VAIK/fhr+hw4ATdIOMwIADwBN0g8zAgAQAE3SEDMCABEATdIRMwIACgBN0gozAgALAE3SCzMCAAwA"
    "TdIMMwIADQBN0g0zAgAGAE3SBjMCAAcATdIHMwIACABN0ggzAgAJAE3SCTMCABoATdICMwIAGwBN0gMzAgACGujRKi5KBIgWBiSI"
    "HhYkCwDt0QsdsgADAO3RAxWyAJAgFCSQGBgkAwAB0gxbDAQKAAHSClssBAoAANIRMSkEAwAA0g0xDQQEAE3SBDMCAAUATdIFMwIA"
    "AAEc2AIKAwCINgYkiA4OJAYA7dEHDbIAAwDt0QM1sgCQEA4kkAgIJAMAAdIEWwwEBAAB0gdbGAQEAADSCTERBAMAANIFMQ0EAgMc"
    "2AIEAwB+Ev6HIywEaP8sLGgAAgAAJyyGff8kJGgAIAAAahCQh/8oKGgAQAAAAAA82AIYAAB+EP6J1QCIvwIAhtIWAwIAAwCF0gJR"
    "AAAWBwZqgQQIaCgGln0BAIC/AgkEACgGCGwDCQYAgQQIaCgGln0BAIC/AgkEACkEBCopBC5sAi4EaAQEhn1qIJK+fhKSiAYAiL8C"
    "GujRKi5KBAIDHNgCExMAAAEc2AITEwASIZK+gAIwfn4S/ojS/4i/AgCF0gIPAAAEGujRKy5SBJ8EBiJ/wIy/AgAI0gIDOQCfCAoi"
    "GAAI0gQDCQQAgFzcGAB/DhCAXNwYAH8KIIBc3BgAfwYwgFzcGAB/AnMPjL+QHDAk/xwyJgAA//+QHjQk/x42JgAA//+QIDgk/yA6"
    "JgAA//+QIjwk/yI+JgAA//8ZAwvRGTMCABgDC9EYMQIAGwML0Rs3AgAaAwvRGjUCAHIPjL+QFEAk/xRCJgAA//+QFkQk/xZGJgAA"
    "//8dAwvRHTsCABwDC9EcOQIAHwML0R8/AgAeAwvRHj0CABgzMBYaNzIWkBhIJP8YSiYAAP//kBpMJP8aTiYAAP//IQML0SFDAgAg"
    "AwvRIEECACMDC9EjRwIAIgML0SJFAgAcOzQWHj82FhgA09EYAWUEcQ+Mv5AMUCT/DFImAAD//5AOVCT/DlYmAAD//yUDC9ElSwIA"
    "JAML0SRJAgAnAwvRJ08CACYDC9EmTQIAIEM4FiJHOhYYANPRGDVuBJAQWCT/EFomAAD//5ASXCT/El4mAAD//ykDC9EpUwIAKAML"
    "0ShRAgArAwvRK1cCACoDC9EqVQIAJEs8FiZPPhYYANPRGDl2BHAPjL+QBGAk/wRiJgAA//+QBmQk/wZmJgAA//8tAwvRLVsCACwD"
    "C9EsWQIALwML0S9fAgAuAwvRLl0CAChTQBYqV0IWGADT0Rg9fgSQCGgk/whqJgAA//+QCmwkMQML0TFjAgAwAwvRMGECADMDC9Ez"
    "ZwIAMgML0TJlAgAsW0QWLl9GFhgA09EYQYYE/woyJgAA//81AwvRNWsCADQDC9E0aQIAMGNIFjJnShYYANPRGEWOBBkDC9EZMwIA"
    "GgML0TZtAgA0a0wWGADT0RhJlgQaMzIWGgDT0RhNZgSANJp8gAIyfoACMH5qIJq+3/6Iv/80MGgAACAAGADI0RgvIQKCMDAewjAw"
    "TPkwMiQVBgYA1v6Cv34M/ocmIiKSggAEIH/AjL8AAwrACAAAABYigIH/BDwmcAAAAI8ALiYhAADSAA55BIAAkL4hLywoEACRvgEA"
    "yNEACQkCEACSvhAAk74QcAR+GACF0hYRAAAAARyzEnAIfh8AhdIkLgIAIgCF0iUuAgAFLIh9hQJAJJ8wMiKCAEYkhgBIIH/AjL8A"
    "AIq/kQCFvwafAJAAmwCPvEYEJgYAAIH/QgYm4P///wIAAdIkAwkEAIUHkBwA6NEDEwgEgAIEfhdASmgmAP/RIwKKBBoACNIMAGEE"
    "gACIvgIDBn4CAwh+AgMKfghKJGifJBQgCEAcaITBB78SFRRogRwMIISAAIWBFBQiBgD/0QY/AgCfFBYiAAD+2QYAAAYKAAjSGgEp"
    "BACAXNwKAH8KFwgQgRCfAJAAmwCPAAB02CYAAB0QAACBAIUAkAACHmgAAMTQBx4CAGoAkoaAAk5+EiCAvhAAiL+FHiAkhh4iJP8g"
    "ICYA/////yIiJsAAAACBHh4gEAD/0RwhRgQQAAHSDwVBBJ8gIiIQAAjSDgBBBACAQNwQAH8nfgD+h/8kJGiAAAAAnyQmIP8cHGiA"
    "AAAAhMEHvxInJGiBHBwghIAAhYEkJCIOAP/RDj8CAJ8kJiIAAP7ZDgAADhIACNIaAUkEAIBc3BIAfxIQ/wCBgAAAAACfAZABmwGP"
    "BAB02CYAACgAAQCBAIUAkAACVGgAAMTQB1QCAGoAkIaAAlJ+ECCAvhAAiL+FVFIkhlRWJP9SUiYA/////1ZWJsAAAACBVFQgKQD/"
    "0RxTrgQqAAHSKgWlBJ9UViIqAAjSDgCpBACAQNwqAH8pfgD+h3/BjL//Ojom/wAAAH/AjL//UFAm/wAAAAj/EIEAAQAAcQ+MvwAA"
    "rNMdTwIAAgyt0wYVCoQI/wCBAAIAAAAcAr+ITExocA+MvwAArNMoUwIAAgyt0w4lCoQCAIW/EACIvoT/gr8QHAO/UgCFvwafAJAA"
    "mwCPvEYMJgYAAIH/Qg4m4P///wYAAdIkAxkEAIUGkBAA6NEHExgEEIUAjx0AHYEdRAxoDgAI0gwAYQQFLIh9F0AiaBIA/9EGAxIC"
    "DQCCv34A/od/wIy//yYmJv8AAACAABC3EBwEv3APjL8AAKzTEysCAAIMrdMGFQqEhCQkaDAAhL8QIhRonxQWIBBADGiEwQe/ChcU"
    "aIEMDCCEgACFgRQUIgYA/9EGPwIAnxQWIgAA/tkGAAAGCgAI0g4BKQQAgFzcCgB/ChcQAIEAnwGQAZsBjwAAdNgSAAATAAEAgQCF"
    "AJAAAihoAADE0AYoAgBqAIiGgAIqfggggL7U/4i/hSgqJIYoMCT/KiomAP////8wMCbAAAAAgSgoIBUA/9EQK2IEFAAB0hQFVQSf"
    "KCoiFAAI0g4AUQQAgEDcFAB/FcP/gr+CAgIkAgIMKB6BAr8AAMTQBAwCAMEBhr6FAIS/BSyIfQQDBpIAapCGECCCvggAiL8GDA5o"
    "CAjo0QcLWASfEBIiCAAI0ggFKQAAgHDcCAJ/AH4C/oeBDBIoAgDE0AQSAgACao6GDiCCvggAiL8GEg5oCgjo0QcLWASfFBYiCgAI"
    "0goFKQAAgHDcCgN/AH4C/oeCDBAoAgDE0AQQAgACaoyGDCCCvggAiL8GEA5oCgjo0QcLWASfFBYiCgAI0goFKQAAgHDcCgR/AH4C"
    "/oeDDA4oAgDE0AQOAgACaoiGCCCCvggAiL8GDhRoCgbo0QoLWASfFBYiCgAI0goFKQAAgHDcCgV/AH4C/ocYgBK/gAGGvkcAhb8C"
    "AMrQgAACAADAoOAAAAAAcA+MvwIgkr4dAIi/fgGcvgAAjNIcAAEAAACN0h0AAgCAAJR9aiCavg4AiL8WnxeQFoKkjhgkJIAZJSWC"
    "HA2XvoACFH4XAhZ+AICg4AAAAAAAgAndCgskCnAPjL8AgKTgAAAAAH4a/ocKBS5+gAIUfgAAgL8XAABoAAAa2AoAAAB+Ev6HgAIA"
    "fn/AjL8AAIq/AABs2AAAAAAewRKBf8CMvxIAmn0ZAIe/BQQXkgDApOAAAAAAECCSvmQAib9+Ev6HDiCQvokAib9+EP6HDCCOvqkA"
    "ib9+Dv6HCCCMvskAib9+DP6HAiCIvgcAiL8WnxeQFoKCjhgCAoAZAwOCgAIAfgCAcNwAAAIAfgj+h34G6oZNAIa/BSyIfQBqgoYC"
    "IIC+DQCIvwAAyNECIQUC/38CsAAA/9ECAQoAkAAEIAAC6NEGC1gEnwACIgAACNIAA1EAAIBp3gACfwB+AP6HgQwAKAAAxNAEAAIA"
    "AGqChgIggL4NAIi/AQDI0QMhBQL/fwKwAQD/0QMDCgCQAgQgAALo0QALWASfAAIiAAAI0gADUQAAgGneAAJ/AH4A/oeCDAAoAADE"
    "0AQAAgAAaoKGAiCAvg0AiL8BAMjRBCEFAv9/ArABAP/RBAMKAJACBCAAAujRAAtYBJ8AAiIAAAjSAANRAACAad4AAn8AfgD+h4MM"
    "ACgAAMTQBAACAABqgIYAIIK+DQCIvwEAyNEFIQUC/38AsAEA/9EFAwIAkAIEIAAA6NEAC1gEnwACIgAACNIAA1EAAIBp3gACfwAA"
    "AIG/AAAA0iEIBQQfhBCOEAAAbAAAhdIFAAIAAAD90RYOAQQAAP/RAD1eBCKHEI4QAABsgAIUfh4AkL6fAAIiDAAI0gAFKQAAgFDc"
    "DAB/ARDBEIEQgAa/FwAAaHAPjL8KAxQC9f+EvwAAyNEKIQUC/38QsAoA/9EKAUIAABDo0QYLWASfAAIiAAAI0gADUQAAgGzcAAp/"
    "AH4S/ocOIJC+d/+IvwkAhdIFEgIAAAD90RYOJQQAAP/RAD1eBCKHDo4OAABsgAIUfh4Ajr6fAAIiDAAI0gAFKQAAgFDcDAB/AQ7B"
    "DoEOgAe/FwAAaHAPjL8KAxQC9f+FvwAAyNEKIQUC/38OsAoA/9EKAToACS0AaJ8AAiIAAAjSAANRAACAbNwACn8AfhD+hwwgjr5X"
    "/4i/CACF0gUQAgAAAP3RFg4hBAAA/9EAPV4EIocMjgwAAGyAAhJ+HgCMvp8AAiIKAAjSAAUpAACAUNwKAH8BDMEMgQyAB78XAABo"
    "cA+MvwkDEgL1/4W/AADI0QkhBQL/fwywCQD/0QkBMgAILQBonwACIgAACNIAA1EAAIBs3AAJfwB+Dv6HCCCMvjf/iL8HAIXSBQ4C"
    "AAAA/dEWDh0EAAD/0QA9XgQihwiOCAAAbIACEH6fAAIiCgAI0gAFKQAAgFDcCgB/AR7BHoEegAe/FwAAaHAPjL8IAxAC9f+FvwAA"
    "yNEIIQUC/38IsAgA/9EIASIABy0AaJ8AAiIAAAjSAANRAACAbNwACH8Afgz+hwIgiL4Y/4m/Hv+CvwAAgL8AAIC/AACAvwAAgL8A"
    "AIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/"
    "AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACA"
    "vwAAgL8AAIC/AACAvwAAgL8AAQbAEAAAAAAAANICEAEEf8CMvwQAiH1qIIK+IQCIv4AABsAIAAAABYEEvxEAhb8AAAbAAAAAAIAC"
    "An4AAwR+nwQGIn/AjL8EAAjSAgUBAACAUNwEAH8DBcEFgQWABr8EBARocA+MvwEHAgL0/4S/AQCCv4ACAn4CAMjRASEFAv9/ALAC"
    "AP/RAQUCAJ8AAiJ/wIy/AAAI0gADCQAAgGzcAAJ/AAAAgb8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/"
    "AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACA"
    "vwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAA"
    "gL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8A"
    "AIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/"
    "AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACA"
    "vwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAA"
    "gL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8A"
    "AIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/"
    "AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACA"
    "vwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAA"
    "gL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8A"
    "AIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/"
    "AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/AACA"
    "vwAAgL8AAIC/AACAvwAAgL8AAIC/AACAvwAAgL8AAIC/BgAAAAAAAAAwDAAAAAAAAAsAAAAAAAAAGAAAAAAAAAAFAAAAAAAAACgN"
    "AAAAAAAACgAAAAAAAAB5AAAAAAAAAPX+/28AAAAAwAwAAAAAAAAEAAAAAAAAAPAMAAAAAAAAAAAAAAAAAAAAAAAAAAAAAABBTUQg"
    "Y2xhbmcgdmVyc2lvbiAyMC4wLjBnaXQgKGh0dHBzOi8vZ2l0aHViLmNvbS9SYWRlb25PcGVuQ29tcHV0ZS9sbHZtLXByb2plY3Qg"
    "cm9jLTcuMS4wIDI1NDI1IDFiMGVhZGE2YjBlZTkzZTJlNjk0YzhjMTQ2ZDIzZmNhOTBiYzExYzUpAExpbmtlcjogQU1EIExMRCAy"
    "MC4wLjAgKC9sb25nZXJfcGF0aG5hbWVfc29fdGhhdF9ycG1zX2Nhbl9zdXBwb3J0X3BhY2thZ2luZ190aGVfZGVidWdfaW5mb19m"
    "b3JfYWxsX29zX3Byb2ZpbGVzL3NyYy9sbHZtLXByb2plY3QvbGx2bSAxYjBlYWRhNmIwZWU5M2UyZTY5NGM4YzE0NmQyM2ZjYTkw"
    "YmMxMWM1KQAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAQAAAAAA8f8AAAAAAAAAAAAAAAAAAAAAKgAAAAAA8f83AAAA"
    "AAAAAAAAAAAAAAAASwAAAAAA8f8AAAAAAAAAAAAAAAAAAAAAbAAAAAAA8f8uAAAAAAAAAAAAAAAAAAAAkgAAAAAA8f8BAAAAAAAA"
    "AAAAAAAAAAAAswAAAAAA8f8AAAAAAAAAAAAAAAAAAAAA3QAAAAAA8f8AAAAAAAAAAAAAAAAAAAAACQEAAAAA8f8AAAAAAAAAAAAA"
    "AAAAAAAALwEAAAAA8f8AAAAAAAAAAAAAAAAAAAAAUwEAAAAA8f8GAAAAAAAAAAAAAAAAAAAAbwEAAAAA8f8AAAAAAAAAAAAAAAAA"
    "AAAAiwEAAAAA8f8GAAAAAAAAAAAAAAAAAAAArAEAAAAA8f8BAAAAAAAAAAAAAAAAAAAAyAEAAAAA8f8AAAAAAAAAAAAAAAAAAAAA"
    "7QEAAAAA8f8AAAAAAAAAAAAAAAAAAAAAFAIAAAAA8f8AAAAAAAAAAAAAAAAAAAAArQIAAAACCADARAAAAAAAAAAAAAAAAAAANQIA"
    "ABIDBwAAHwAAAAAAAEgQAAAAAAAATQIAABEDBgDADQAAAAAAAEAAAAAAAAAAaAIAABIDBwAAMAAAAAAAAKgAAAAAAAAAewIAABED"
    "BgAADgAAAAAAAEAAAAAAAAAAkQIAABEACgAwVQAAAAAAAAEAAAAAAAAAAC5ub3RlAC5keW5zeW0ALmdudS5oYXNoAC5oYXNoAC5k"
    "eW5zdHIALnJvZGF0YQAudGV4dAAuZHluYW1pYwAucmVscm9fcGFkZGluZwAuYnNzAC5BTURHUFUuZ3ByX21heGltdW1zAC5jb21t"
    "ZW50AC5zeW10YWIALnNoc3RydGFiAC5zdHJ0YWIAAGZ1c2VkX214ZnA0X2dlbW1fa2VybmVsLnByaXZhdGVfc2VnX3NpemUAZnVz"
    "ZWRfbXhmcDRfZ2VtbV9rZXJuZWwubnVtX3ZncHIAZnVzZWRfbXhmcDRfZ2VtbV9rZXJuZWwubnVtX2FncHIAZnVzZWRfbXhmcDRf"
    "Z2VtbV9rZXJuZWwubnVtYmVyZWRfc2dwcgBmdXNlZF9teGZwNF9nZW1tX2tlcm5lbC51c2VzX3ZjYwBmdXNlZF9teGZwNF9nZW1t"
    "X2tlcm5lbC51c2VzX2ZsYXRfc2NyYXRjaABmdXNlZF9teGZwNF9nZW1tX2tlcm5lbC5oYXNfZHluX3NpemVkX3N0YWNrAGZ1c2Vk"
    "X214ZnA0X2dlbW1fa2VybmVsLmhhc19yZWN1cnNpb24AcmVkdWNlX2JmMTZfa2VybmVsLnByaXZhdGVfc2VnX3NpemUAcmVkdWNl"
    "X2JmMTZfa2VybmVsLm51bV92Z3ByAHJlZHVjZV9iZjE2X2tlcm5lbC5udW1fYWdwcgByZWR1Y2VfYmYxNl9rZXJuZWwubnVtYmVy"
    "ZWRfc2dwcgByZWR1Y2VfYmYxNl9rZXJuZWwudXNlc192Y2MAcmVkdWNlX2JmMTZfa2VybmVsLnVzZXNfZmxhdF9zY3JhdGNoAHJl"
    "ZHVjZV9iZjE2X2tlcm5lbC5oYXNfZHluX3NpemVkX3N0YWNrAHJlZHVjZV9iZjE2X2tlcm5lbC5oYXNfcmVjdXJzaW9uAGZ1c2Vk"
    "X214ZnA0X2dlbW1fa2VybmVsAGZ1c2VkX214ZnA0X2dlbW1fa2VybmVsLmtkAHJlZHVjZV9iZjE2X2tlcm5lbAByZWR1Y2VfYmYx"
    "Nl9rZXJuZWwua2QAX19oaXBfY3VpZF85OTk1NmI5MWU0ODIyNWU5AF9EWU5BTUlDAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAABAAAABwAAAAIAAAAAAAAAOAIAAAAAAAA4AgAAAAAAAPgJAAAA"
    "AAAAAAAAAAAAAAAEAAAAAAAAAAAAAAAAAAAABwAAAAsAAAACAAAAAAAAADAMAAAAAAAAMAwAAAAAAACQAAAAAAAAAAUAAAABAAAA"
    "CAAAAAAAAAAYAAAAAAAAAA8AAAD2//9vAgAAAAAAAADADAAAAAAAAMAMAAAAAAAAMAAAAAAAAAACAAAAAAAAAAgAAAAAAAAAAAAA"
    "AAAAAAAZAAAABQAAAAIAAAAAAAAA8AwAAAAAAADwDAAAAAAAADgAAAAAAAAAAgAAAAAAAAAEAAAAAAAAAAQAAAAAAAAAHwAAAAMA"
    "AAACAAAAAAAAACgNAAAAAAAAKA0AAAAAAAB5AAAAAAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAAACcAAAABAAAAAgAAAAAAAADA"
    "DQAAAAAAAMANAAAAAAAAgAAAAAAAAAAAAAAAAAAAAEAAAAAAAAAAAAAAAAAAAAAvAAAAAQAAAAYAAAAAAAAAAB8AAAAAAAAADwAA"
    "AAAAAMAVAAAAAAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAANQAAAAYAAAADAAAAAAAAAMBEAAAAAAAAwCQAAAAAAABwAAAAAAAA"
    "AAUAAAAAAAAACAAAAAAAAAAQAAAAAAAAAD4AAAAIAAAAAwAAAAAAAAAwRQAAAAAAADAlAAAAAAAA0AoAAAAAAAAAAAAAAAAAAAEA"
    "AAAAAAAAAAAAAAAAAABNAAAACAAAAAMAAAAAAAAAMFUAAAAAAAAwJQAAAAAAAAEAAAAAAAAAAAAAAAAAAAABAAAAAAAAAAAAAAAA"
    "AAAAUgAAAAEAAAAAAAAAAAAAAAAAAAAAAAAAMCUAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAAAGcAAAABAAAA"
    "MAAAAAAAAAAAAAAAAAAAADAlAAAAAAAAOQEAAAAAAAAAAAAAAAAAAAEAAAAAAAAAAQAAAAAAAABwAAAAAgAAAAAAAAAAAAAAAAAA"
    "AAAAAABwJgAAAAAAACgCAAAAAAAADwAAABIAAAAIAAAAAAAAABgAAAAAAAAAeAAAAAMAAAAAAAAAAAAAAAAAAAAAAAAAmCgAAAAA"
    "AACKAAAAAAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAAAIIAAAADAAAAAAAAAAAAAAAAAAAAAAAAACIpAAAAAAAAtgIAAAAAAAAA"
    "AAAAAAAAAAEAAAAAAAAAAAAAAAAAAAABGwM7RAAAAAcAAABIEQAAYAAAAIgRAACAAAAA2BEAAKAAAAAIEwAAxAAAAJgTAADoAAAA"
    "OBQAAAgBAABoFAAAIAEAAAAAAAAUAAAAAAAAAAF6UgABeBABGwwHCJABAAAcAAAAHAAAAOAQAAA4AAAAAEEOEIYCQw0GcwwHCAAA"
    "ABwAAAA8AAAAABEAAEoAAAAAQQ4QhgJDDQYCRQwHCAAAIAAAAFwAAAAwEQAAJAEAAABHDtABAwwBDtgBRA7gAUwOCAAAIAAAAIAA"
    "AAA8EgAAiwAAAABEDoABAnYOiAFEDpABTA4IAAAAHAAAAKQAAACoEgAAlAAAAABBDhBEDjCDAgKJDhBBDggUAAAAxAAAACgTAAAi"
    "AAAAAE0OEFQOCAAUAAAA3AAAAEATAAASAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAVUiJ5fYFnSYAAAF0AusnxgWSJgAAAUiL"
    "BeMVAABIg/gAdBNIjT3W7v//SI01fyYAAOg6AwAAXcMPH4QAAAAAAFVIieX2BaUmAAABdALrOcYFmiYAAAFIiwWrFQAASIP4AHQM"
    "SIs9viUAAOgRAwAASIsFmhUAAEiD+AB0DEiNPX3u///oCAMAAF3DzMzMzMzMSIHsyAAAAEiJfCRYSIl0JFBIiVQkSEiJTCRARIlE"
    "JAxEiUwkCEiNRCRYSIlEJGBIjUQkUEiJRCRoSI1EJEhIiUQkcEiNRCRASIlEJHhIjUQkDEiJhCSAAAAASI1EJAhIiYQkiAAAAEiN"
    "hCTQAAAASImEJJAAAABIjYQk2AAAAEiJhCSYAAAASI2EJOAAAABIiYQkoAAAAEiNhCToAAAASImEJKgAAABIjYQk8AAAAEiJhCSw"
    "AAAASI2EJPgAAABIiYQkuAAAAEiNhCQAAQAASImEJMAAAABIjXwkMEiNdCQgSI1UJBhIjUwkEOggAgAASIt0JDCLVCQ4SItMJCBE"
    "i0QkKEiLPX4UAABMjUwkYP90JBD/dCQg6AQCAABIgcTYAAAAw2ZmZi4PH4QAAAAAAEiD7HhIiXwkSEiJdCRAiVQkDIlMJAhIjUQk"
    "SEiJRCRQSI1EJEBIiUQkWEiNRCQMSIlEJGBIjUQkCEiJRCRoSI18JDBIjXQkIEiNVCQYSI1MJBDoiQEAAEiLdCQwi1QkOEiLTCQg"
    "RItEJChIiz3vEwAATI1MJFD/dCQQ/3QkIOhtAQAASIHEiAAAAMMPH0QAAFNIg+wgSIsdnCQAAEiF23UWSI09yCMAAOhTAQAASInD"
    "SIkFgSQAAA9XwA8RRCQQDxEEJEiLNY4TAABIjQ2qpP//SInfSInKQbj/////RTHJ6CsBAAAPV8APEUQkEA8RBCRIizVoEwAASI0N"
    "aaT//0iJ30iJykG4/////0Uxyej9AAAASI09FgAAAEiDxCBb6TwAAABmZmYuDx+EAAAAAABIiz0BJAAASIX/dBVQ6N4AAABIxwXr"
    "IwAAAAAAAEiDxAjDzMzMzMzMzMzMzMzMzMzzDx76SIsV/SIAADH26b4AAADMzPMPHvpIg+wISIsFtRIAAEiFwHQC/9BIg8QIwwDz"
    "Dx76SIPsCEiDxAjDAAAA/zXqIgAA/yXsIgAADx9AAP8l6iIAAGgAAAAA6eD/////JeIiAABoAQAAAOnQ/////yXaIgAAaAIAAADp"
    "wP////8l0iIAAGgDAAAA6bD/////JcoiAABoBAAAAOmg/////yXCIgAAaAUAAADpkP////8luiIAAGgGAAAA6YD/////JbIiAABo"
    "BwAAAOlw/////yWqIgAAaAgAAADpYP///wAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAdAAAAAAAAAJ8B"
    "AAAAAAAAAQAAAAAAAABaAQAAAAAAAAEAAAAAAAAAswEAAAAAAAABAAAAAAAAAMIBAAAAAAAAAQAAAAAAAABzAQAAAAAAAAEAAAAA"
    "AAAAiQEAAAAAAAAHAAAAAAAAADgGAAAAAAAACAAAAAAAAAA4AQAAAAAAAAkAAAAAAAAAGAAAAAAAAAD5//9vAAAAAAUAAAAAAAAA"
    "FwAAAAAAAABwBwAAAAAAAAIAAAAAAAAA2AAAAAAAAAADAAAAAAAAAGiHAAAAAAAAFAAAAAAAAAAHAAAAAAAAAAYAAAAAAAAAOAIA"
    "AAAAAAALAAAAAAAAABgAAAAAAAAABQAAAAAAAABoBAAAAAAAAAoAAAAAAAAAzAEAAAAAAAD1/v9vAAAAADgEAAAAAAAAGQAAAAAA"
    "AAAgdQAAAAAAABsAAAAAAAAAEAAAAAAAAAAaAAAAAAAAADB1AAAAAAAAHAAAAAAAAAAIAAAAAAAAAAwAAAAAAAAAVGQAAAAAAAAN"
    "AAAAAAAAAHBkAAAAAAAA8P//bwAAAAC4AwAAAAAAAP7//28AAAAA2AMAAAAAAAD///9vAAAAAAMAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAARlBJSAEAAAAAAAAAAAAA"
    "AAAAAAAAAAAASHUAAAAAAAAAAAAAAAAAAAAAAAAAAAAAlmQAAAAAAACmZAAAAAAAALZkAAAAAAAAxmQAAAAAAADWZAAAAAAAAOZk"
    "AAAAAAAA9mQAAAAAAAAGZQAAAAAAABZlAAAAAAAAAEFNRCBjbGFuZyB2ZXJzaW9uIDIwLjAuMGdpdCAoaHR0cHM6Ly9naXRodWIu"
    "Y29tL1JhZGVvbk9wZW5Db21wdXRlL2xsdm0tcHJvamVjdCByb2MtNy4xLjAgMjU0MjUgMWIwZWFkYTZiMGVlOTNlMmU2OTRjOGMx"
    "NDZkMjNmY2E5MGJjMTFjNSkATGlua2VyOiBBTUQgTExEIDIwLjAuMCAoL2xvbmdlcl9wYXRobmFtZV9zb190aGF0X3JwbXNfY2Fu"
    "X3N1cHBvcnRfcGFja2FnaW5nX3RoZV9kZWJ1Z19pbmZvX2Zvcl9hbGxfb3NfcHJvZmlsZXMvc3JjL2xsdm0tcHJvamVjdC9sbHZt"
    "IDFiMGVhZGE2YjBlZTkzZTJlNjk0YzhjMTQ2ZDIzZmNhOTBiYzExYzUpAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAB"
    "AAAABADx/wAAAAAAAAAAAAAAAAAAAAAMAAAAAgAMACBhAAAAAAAAOAAAAAAAAAAWAAAAAQAZAMiHAAAAAAAAAQAAAAAAAAAuAAAA"
    "AQALACBQAAAAAAAAAAAAAAAAAABAAAAAAQAZANCHAAAAAAAAQAAAAAAAAABTAAAAAgAMAGBhAAAAAAAASgAAAAAAAABdAAAAAQAZ"
    "ABCIAAAAAAAAAQAAAAAAAABzAAAAAQAQACB1AAAAAAAACAAAAAAAAAB6AAAAAQARADB1AAAAAAAACAAAAAAAAABNAQAAAQIWAEiH"
    "AAAAAAAACAAAAAAAAACBAAAABADx/wAAAAAAAAAAAAAAAAAAAACQAAAAAgAMAHBjAAAAAAAAlAAAAAAAAACiAAAAAQAZABiIAAAA"
    "AAAACAAAAAAAAADHAAAAAQAXAFCHAAAAAAAAGAAAAAAAAADcAAAAAgAMABBkAAAAAAAAIgAAAAAAAADuAAAABADx/wAAAAAAAAAA"
    "AAAAAAAAAACGAgAAAQILACBQAAAAAAAABAAAAAAAAAAGAQAAAgINAFRkAAAAAAAAAAAAAAAAAAAiAQAAAgIOAHBkAAAAAAAAAAAA"
    "AAAAAAAMAQAAAAIYAGiHAAAAAAAAAAAAAAAAAACcAgAAAAITAEh1AAAAAAAAAAAAAAAAAAA9AgAAAgIMAEBkAAAAAAAAEgAAAAAA"
    "AAD3AAAAIAAAAAAAAAAAAAAAAAAAAAAAAAAoAQAAIgAAAAAAAAAAAAAAAAAAAAAAAAA+AQAAIgAAAAAAAAAAAAAAAAAAAAAAAABa"
    "AQAAIgAAAAAAAAAAAAAAAAAAAAAAAAByAQAAEgAMALBhAAAAAAAAJAEAAAAAAACZAQAAEgAAAAAAAAAAAAAAAAAAAAAAAACzAQAA"
    "EQASADh1AAAAAAAACAAAAAAAAADLAQAAEgAAAAAAAAAAAAAAAAAAAAAAAADbAQAAEgAMAOBiAAAAAAAAiwAAAAAAAAD9AQAAEQAS"
    "AEB1AAAAAAAACAAAAAAAAAAQAgAAEgAAAAAAAAAAAAAAAAAAAAAAAAAnAgAAEgAAAAAAAAAAAAAAAAAAAAAAAABEAgAAEgAAAAAA"
    "AAAAAAAAAAAAAAAAAABdAgAAEQAZACCIAAAAAAAAAQAAAAAAAAB5AgAAEgAAAAAAAAAAAAAAAAAAAAAAAAAALmR5bnN5bQAuZ251"
    "LnZlcnNpb24ALmdudS52ZXJzaW9uX3IALmdudS5oYXNoAC5keW5zdHIALnJlbGEuZHluAC5yZWxhLnBsdAAucm9kYXRhAC5oaXBf"
    "ZmF0YmluAC5laF9mcmFtZV9oZHIALmVoX2ZyYW1lAC50ZXh0AC5pbml0AC5maW5pAC5wbHQALmluaXRfYXJyYXkALmZpbmlfYXJy"
    "YXkALmRhdGEucmVsLnJvAC5keW5hbWljAC5nb3QALnJlbHJvX3BhZGRpbmcALmRhdGEALmhpcEZhdEJpblNlZ21lbnQALmdvdC5w"
    "bHQALmJzcwAuY29tbWVudAAuc3ltdGFiAC5zaHN0cnRhYgAuc3RydGFiAABjcnRiZWdpbi5jAF9fZG9faW5pdABfX2RvX2luaXQu"
    "X19pbml0aWFsaXplZABfX0VIX0ZSQU1FX0xJU1RfXwBfX2RvX2luaXQuX19vYmplY3QAX19kb19maW5pAF9fZG9fZmluaS5fX2Zp"
    "bmFsaXplZABfX2luaXQAX19maW5pAHYxNV9rZXJuZWwuaGlwAF9faGlwX21vZHVsZV9jdG9yAF9faGlwX2dwdWJpbl9oYW5kbGVf"
    "OTk5NTZiOTFlNDgyMjVlOQBfX2hpcF9mYXRiaW5fd3JhcHBlcgBfX2hpcF9tb2R1bGVfZHRvcgBjcnRlbmQuYwBfX2dtb25fc3Rh"
    "cnRfXwBfaW5pdABfR0xPQkFMX09GRlNFVF9UQUJMRV8AX2ZpbmkAX19yZWdpc3Rlcl9mcmFtZV9pbmZvAF9fY3hhX2ZpbmFsaXpl"
    "AF9fZHNvX2hhbmRsZQBfX2RlcmVnaXN0ZXJfZnJhbWVfaW5mbwBfX2RldmljZV9zdHViX19mdXNlZF9teGZwNF9nZW1tX2tlcm5l"
    "bABfX2hpcFBvcENhbGxDb25maWd1cmF0aW9uAGZ1c2VkX214ZnA0X2dlbW1fa2VybmVsAGhpcExhdW5jaEtlcm5lbABfX2Rldmlj"
    "ZV9zdHViX19yZWR1Y2VfYmYxNl9rZXJuZWwAcmVkdWNlX2JmMTZfa2VybmVsAF9faGlwUmVnaXN0ZXJGYXRCaW5hcnkAX19oaXBS"
    "ZWdpc3RlckZ1bmN0aW9uAGF0ZXhpdABfX2hpcFVucmVnaXN0ZXJGYXRCaW5hcnkAX19oaXBfY3VpZF85OTk1NmI5MWU0ODIyNWU5"
    "AF9fY3hhX2F0ZXhpdABfX0VIX0ZSQU1FX0xJU1RfRU5EX18AX0RZTkFNSUMAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA"
    "AAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAQAAAAsAAAACAAAAAAAAADgCAAAAAAAAOAIAAAAAAACAAQAAAAAA"
    "AAUAAAABAAAACAAAAAAAAAAYAAAAAAAAAAkAAAD///9vAgAAAAAAAAC4AwAAAAAAALgDAAAAAAAAIAAAAAAAAAABAAAAAAAAAAIA"
    "AAAAAAAAAgAAAAAAAAAWAAAA/v//bwIAAAAAAAAA2AMAAAAAAADYAwAAAAAAAGAAAAAAAAAABQAAAAMAAAAEAAAAAAAAAAAAAAAA"
    "AAAAJQAAAPb//28CAAAAAAAAADgEAAAAAAAAOAQAAAAAAAAwAAAAAAAAAAEAAAAAAAAACAAAAAAAAAAAAAAAAAAAAC8AAAADAAAA"
    "AgAAAAAAAABoBAAAAAAAAGgEAAAAAAAAzAEAAAAAAAAAAAAAAAAAAAEAAAAAAAAAAAAAAAAAAAA3AAAABAAAAAIAAAAAAAAAOAYA"
    "AAAAAAA4BgAAAAAAADgBAAAAAAAAAQAAAAAAAAAIAAAAAAAAABgAAAAAAAAAQQAAAAQAAABCAAAAAAAAAHAHAAAAAAAAcAcAAAAA"
    "AADYAAAAAAAAAAEAAAAYAAAACAAAAAAAAAAYAAAAAAAAAEsAAAABAAAAMgAAAAAAAABICAAAAAAAAEgIAAAAAAAAKwAAAAAAAAAA"
    "AAAAAAAAAAEAAAAAAAAAAQAAAAAAAABTAAAAAQAAAAIAAAAAAAAAABAAAAAAAAAAEAAAAAAAANg/AAAAAAAAAAAAAAAAAAAAEAAA"
    "AAAAAAAAAAAAAAAAXwAAAAEAAAACAAAAAAAAANhPAAAAAAAA2E8AAAAAAABEAAAAAAAAAAAAAAAAAAAABAAAAAAAAAAAAAAAAAAA"
    "AG0AAAABAAAAAgAAAAAAAAAgUAAAAAAAACBQAAAAAAAA9AAAAAAAAAAAAAAAAAAAAAgAAAAAAAAAAAAAAAAAAAB3AAAAAQAAAAYA"
    "AAAAAAAAIGEAAAAAAAAgUQAAAAAAADQDAAAAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAAAAfQAAAAEAAAAGAAAAAAAAAFRkAAAA"
    "AAAAVFQAAAAAAAAbAAAAAAAAAAAAAAAAAAAABAAAAAAAAAAAAAAAAAAAAIMAAAABAAAABgAAAAAAAABwZAAAAAAAAHBUAAAAAAAA"
    "DQAAAAAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAAACJAAAAAQAAAAYAAAAAAAAAgGQAAAAAAACAVAAAAAAAAKAAAAAAAAAAAAAA"
    "AAAAAAAQAAAAAAAAAAAAAAAAAAAAjgAAAA4AAAADAAAAAAAAACB1AAAAAAAAIFUAAAAAAAAQAAAAAAAAAAAAAAAAAAAACAAAAAAA"
    "AAAAAAAAAAAAAJoAAAAPAAAAAwAAAAAAAAAwdQAAAAAAADBVAAAAAAAACAAAAAAAAAAAAAAAAAAAAAgAAAAAAAAAAAAAAAAAAACm"
    "AAAAAQAAAAMAAAAAAAAAOHUAAAAAAAA4VQAAAAAAABAAAAAAAAAAAAAAAAAAAAAIAAAAAAAAAAAAAAAAAAAAswAAAAYAAAADAAAA"
    "AAAAAEh1AAAAAAAASFUAAAAAAADQAQAAAAAAAAUAAAAAAAAACAAAAAAAAAAQAAAAAAAAALwAAAABAAAAAwAAAAAAAAAYdwAAAAAA"
    "ABhXAAAAAAAAMAAAAAAAAAAAAAAAAAAAAAgAAAAAAAAAAAAAAAAAAADBAAAACAAAAAMAAAAAAAAASHcAAAAAAABIVwAAAAAAALgI"
    "AAAAAAAAAAAAAAAAAAABAAAAAAAAAAAAAAAAAAAA0AAAAAEAAAADAAAAAAAAAEiHAAAAAAAASFcAAAAAAAAIAAAAAAAAAAAAAAAA"
    "AAAACAAAAAAAAAAAAAAAAAAAANYAAAABAAAAAwAAAAAAAABQhwAAAAAAAFBXAAAAAAAAGAAAAAAAAAAAAAAAAAAAAAgAAAAAAAAA"
    "AAAAAAAAAADoAAAAAQAAAAMAAAAAAAAAaIcAAAAAAABoVwAAAAAAAGAAAAAAAAAAAAAAAAAAAAAIAAAAAAAAAAAAAAAAAAAA8QAA"
    "AAgAAAADAAAAAAAAAMiHAAAAAAAAyFcAAAAAAABZAAAAAAAAAAAAAAAAAAAACAAAAAAAAAAAAAAAAAAAAPYAAAABAAAAMAAAAAAA"
    "AAAAAAAAAAAAAMhXAAAAAAAAOQEAAAAAAAAAAAAAAAAAAAEAAAAAAAAAAQAAAAAAAAD/AAAAAgAAAAAAAAAAAAAAAAAAAAAAAAAI"
    "WQAAAAAAAJADAAAAAAAAHQAAABcAAAAIAAAAAAAAABgAAAAAAAAABwEAAAMAAAAAAAAAAAAAAAAAAAAAAAAAmFwAAAAAAAAZAQAA"
    "AAAAAAAAAAAAAAAAAQAAAAAAAAAAAAAAAAAAABEBAAADAAAAAAAAAAAAAAAAAAAAAAAAALFdAAAAAAAApQIAAAAAAAAAAAAAAAAA"
    "AAEAAAAAAAAAAAAAAAAAAAA="
)
_CO_BYTES = base64.b64decode(_CO_B64)

# ═══ HIP init (fallback path) ═══
_hip = ctypes.cdll.LoadLibrary("libamdhip64.so")
_hip.hipModuleLaunchKernel.argtypes = [
    ctypes.c_void_p, ctypes.c_uint, ctypes.c_uint, ctypes.c_uint,
    ctypes.c_uint, ctypes.c_uint, ctypes.c_uint,
    ctypes.c_uint, ctypes.c_void_p, ctypes.c_void_p,
    ctypes.POINTER(ctypes.c_void_p)]
_hip.hipModuleLaunchKernel.restype = ctypes.c_int

_fn = None
_fn_reduce = None

def _init_hip():
    global _fn, _fn_reduce
    if _CO_BYTES is None:
        return
    try:
        td = tempfile.mkdtemp()
        p = os.path.join(td, "k.co")
        with open(p, "wb") as f:
            f.write(_CO_BYTES)
        mod = ctypes.c_void_p()
        if _hip.hipModuleLoad(ctypes.byref(mod), p.encode()) == 0:
            fn = ctypes.c_void_p()
            if _hip.hipModuleGetFunction(ctypes.byref(fn), mod, b"fused_mxfp4_gemm_kernel") == 0:
                _fn = fn
            fn2 = ctypes.c_void_p()
            if _hip.hipModuleGetFunction(ctypes.byref(fn2), mod, b"reduce_bf16_kernel") == 0:
                _fn_reduce = fn2
    except:
        pass

_init_hip()

# ═══ HSA AQL init ═══
_hsa = None
_aql_ready = False
_aql_queue = None
_aql_qsize = 0
_aql_qbase = 0
_aql_doorbell = None
_aql_signal = None
_aql_ko = 0       # kernel_object (main kernel)
_aql_priv = 0     # private_segment_size
_aql_rko = 0      # kernel_object (reduce kernel)
_aql_rpriv = 0
_aql_ka_ptr = None # kernarg device pointer (main, 256 bytes)
_aql_rka_ptr = None # kernarg device pointer (reduce, 256 bytes)
_aql_log = []

def _init_aql():
    global _hsa, _aql_ready, _aql_queue, _aql_qsize, _aql_qbase, _aql_doorbell
    global _aql_signal, _aql_ko, _aql_priv, _aql_ka_ptr, _aql_rko, _aql_rpriv, _aql_rka_ptr

    if _CO_BYTES is None:
        _aql_log.append("No .co bytes")
        return

    try:
        h = ctypes.cdll.LoadLibrary("libhsa-runtime64.so")
        _hsa = h
    except:
        _aql_log.append("Cannot load libhsa-runtime64.so")
        return

    # Find GPU agent
    agents = []
    def acb(agent, data):
        agents.append(_ag(agent.handle))
        return 0
    h.hsa_iterate_agents(_AGENT_CB(acb), None)
    gpu = None
    for a in agents:
        dt = ctypes.c_uint32()
        h.hsa_agent_get_info(a, _HSA_AGENT_DEVICE, ctypes.byref(dt))
        if dt.value == _HSA_DEV_GPU:
            gpu = a; break
    if gpu is None:
        _aql_log.append("No GPU agent"); return

    # Create queue
    qp = ctypes.POINTER(_q)()
    err = h.hsa_queue_create(gpu, ctypes.c_uint32(1024), ctypes.c_uint32(_HSA_Q_SINGLE),
                             None, None, ctypes.c_uint32(0), ctypes.c_uint32(0), ctypes.byref(qp))
    if err != 0:
        _aql_log.append(f"Queue create: {err}"); return
    _aql_queue = qp
    qq = qp.contents
    _aql_qsize = qq.size
    _aql_qbase = qq.base_address
    _aql_doorbell = qq.doorbell_signal
    _aql_log.append(f"Queue: size={_aql_qsize} base=0x{_aql_qbase:x}")

    # Create signal
    sig = _sig()
    err = h.hsa_signal_create(ctypes.c_int64(1), ctypes.c_uint32(0), None, ctypes.byref(sig))
    if err != 0:
        _aql_log.append(f"Signal create: {err}"); return
    _aql_signal = sig

    # Find kernarg region
    regions = []
    def rcb(rh, data):
        regions.append(rh)
        return 0
    h.hsa_agent_iterate_regions(gpu, _REGION_CB(rcb), None)
    ka_region = None
    for r in regions:
        seg = ctypes.c_uint32()
        h.hsa_region_get_info(ctypes.c_uint64(r), _HSA_REG_SEGMENT, ctypes.byref(seg))
        if seg.value == _HSA_SEG_GLOBAL:
            fl = ctypes.c_uint32()
            h.hsa_region_get_info(ctypes.c_uint64(r), _HSA_REG_GFLAGS, ctypes.byref(fl))
            if fl.value & _HSA_FLAG_KERNARG:
                ka_region = r
    if ka_region is None:
        _aql_log.append("No kernarg region"); return

    # Alloc kernarg (256 bytes, reusable)
    ka = ctypes.c_void_p()
    err = h.hsa_memory_allocate(ctypes.c_uint64(ka_region), ctypes.c_size_t(256), ctypes.byref(ka))
    if err != 0:
        _aql_log.append(f"Kernarg alloc: {err}"); return
    _aql_ka_ptr = ka
    _aql_log.append(f"Kernarg: 0x{ka.value:x}")

    # Load kernel via HSA
    co_buf = (ctypes.c_char * len(_CO_BYTES))(*_CO_BYTES)

    exe = ctypes.c_uint64()
    err = h.hsa_executable_create_alt(ctypes.c_uint32(_HSA_PROFILE_FULL),
                                      ctypes.c_uint32(_HSA_ROUND_DEFAULT),
                                      None, ctypes.byref(exe))
    if err != 0:
        _aql_log.append(f"Exe create: {err}"); return

    reader = ctypes.c_uint64()
    err = h.hsa_code_object_reader_create_from_memory(co_buf, ctypes.c_size_t(len(_CO_BYTES)),
                                                       ctypes.byref(reader))
    if err != 0:
        _aql_log.append(f"Reader create: {err}"); return

    err = h.hsa_executable_load_agent_code_object(ctypes.c_uint64(exe.value), gpu,
                                                   ctypes.c_uint64(reader.value), None, None)
    if err != 0:
        _aql_log.append(f"Load code obj: {err}"); return

    err = h.hsa_executable_freeze(ctypes.c_uint64(exe.value), None)
    if err != 0:
        _aql_log.append(f"Freeze: {err}"); return

    # Get symbol
    sym = ctypes.c_uint64()
    # Try with .kd suffix first (AMD convention)
    err = h.hsa_executable_get_symbol_by_name(ctypes.c_uint64(exe.value),
            b"fused_mxfp4_gemm_kernel.kd", ctypes.byref(gpu), ctypes.byref(sym))
    if err != 0:
        err = h.hsa_executable_get_symbol_by_name(ctypes.c_uint64(exe.value),
                b"fused_mxfp4_gemm_kernel", ctypes.byref(gpu), ctypes.byref(sym))
        if err != 0:
            _aql_log.append(f"Symbol lookup: {err}"); return

    ko = ctypes.c_uint64()
    h.hsa_executable_symbol_get_info(ctypes.c_uint64(sym.value), _HSA_SYM_KERNEL_OBJ, ctypes.byref(ko))
    priv = ctypes.c_uint32()
    h.hsa_executable_symbol_get_info(ctypes.c_uint64(sym.value), _HSA_SYM_PRIV_SEG, ctypes.byref(priv))

    _aql_ko = ko.value
    _aql_priv = priv.value
    _aql_log.append(f"Kernel: obj=0x{_aql_ko:x} priv={_aql_priv}")

    # Load reduce kernel symbol
    rsym = ctypes.c_uint64()
    err = h.hsa_executable_get_symbol_by_name(ctypes.c_uint64(exe.value),
            b"reduce_bf16_kernel.kd", ctypes.byref(gpu), ctypes.byref(rsym))
    if err != 0:
        err = h.hsa_executable_get_symbol_by_name(ctypes.c_uint64(exe.value),
                b"reduce_bf16_kernel", ctypes.byref(gpu), ctypes.byref(rsym))
    if err == 0:
        rko = ctypes.c_uint64()
        h.hsa_executable_symbol_get_info(ctypes.c_uint64(rsym.value), _HSA_SYM_KERNEL_OBJ, ctypes.byref(rko))
        rpriv = ctypes.c_uint32()
        h.hsa_executable_symbol_get_info(ctypes.c_uint64(rsym.value), _HSA_SYM_PRIV_SEG, ctypes.byref(rpriv))
        _aql_rko = rko.value
        _aql_rpriv = rpriv.value
        _aql_log.append(f"Reduce: obj=0x{_aql_rko:x} priv={_aql_rpriv}")
        # Alloc separate kernarg for reduce
        rka = ctypes.c_void_p()
        err2 = h.hsa_memory_allocate(ctypes.c_uint64(ka_region), ctypes.c_size_t(256), ctypes.byref(rka))
        if err2 == 0:
            _aql_rka_ptr = rka
    else:
        _aql_log.append(f"Reduce symbol not found: {err}")

    _aql_ready = True
    _aql_log.append("AQL READY")

    # Keep references alive
    global _aql_gc
    _aql_gc = (co_buf, exe, reader, gpu)

_init_aql()
# Print AQL init log
for line in _aql_log:
    print(f"[AQL] {line}", file=sys.stderr)

# ═══ Dispatch ═══
_cache = {}
_gc = {}

def _get_split_k(m, n, k):
    # With AQL dispatch (~1.5 µs), sk=1 is faster than sk=8 + reduce (~7 µs HIP overhead)
    # Even for case 2 (17 tiles, 6.6% CU util), the saved dispatch > lost parallelism
    return 1

def _setup(key, data):
    m, n, k = key
    sa = data[0].stride(0)
    sbq = data[2].view(torch.uint8).stride(0)
    ksp = ((k // 32 + 7) // 8) * 8
    sk = _get_split_k(m, n, k)
    k_chunk = k // sk
    ntn = (n + 127) // 128; ntm = (m + 15) // 16; gx = ntm * ntn
    local_K = k_chunk
    lds = 16 * (local_K // 2) + 16 * (local_K // 32) + 4
    C = torch.empty(m, n, dtype=torch.bfloat16, device="cuda")
    c_ptr = C.data_ptr()

    w_ptr = 0; td_ptr = 0; W = None; td_obj = None; need_reduce = False
    if sk > 1:
        W = torch.empty(sk * m * n, dtype=torch.float32, device="cuda")
        w_ptr = W.data_ptr(); td_ptr = 0; need_reduce = True

    # HIP kernarg
    buf = (ctypes.c_char * 76)()
    struct.pack_into("<QQQQ6iQQi", buf, 0, 0, 0, 0, c_ptr,
                     m, n, k, sa, sbq, ksp, w_ptr, td_ptr, k_chunk)
    buf_ptr = ctypes.cast(buf, ctypes.c_void_p)
    asz = ctypes.c_size_t(76)
    cfg = (ctypes.c_void_p * 5)(
        ctypes.c_void_p(1), buf_ptr,
        ctypes.c_void_p(2), ctypes.cast(ctypes.byref(asz), ctypes.c_void_p),
        ctypes.c_void_p(3))
    cfg_ptr = ctypes.cast(cfg, ctypes.POINTER(ctypes.c_void_p))

    r_cfg_ptr = None; r_gx = 0; _r_keep = []
    if need_reduce and _fn_reduce:
        num_elems = m * n; r_gx = (num_elems + 255) // 256
        r_buf = (ctypes.c_char * 24)()
        struct.pack_into("<QQii", r_buf, 0, w_ptr, c_ptr, num_elems, sk)
        r_bp = ctypes.cast(r_buf, ctypes.c_void_p)
        r_asz = ctypes.c_size_t(24)
        r_cfg = (ctypes.c_void_p * 5)(
            ctypes.c_void_p(1), r_bp,
            ctypes.c_void_p(2), ctypes.cast(ctypes.byref(r_asz), ctypes.c_void_p),
            ctypes.c_void_p(3))
        r_cfg_ptr = ctypes.cast(r_cfg, ctypes.POINTER(ctypes.c_void_p))
        _r_keep = [r_buf, r_bp, r_asz, r_cfg]

    use_aql = _aql_ready and not need_reduce  # AQL only for non-split-K
    # Pre-pack reduce kernarg bytes
    r_ka_bytes = None
    if need_reduce:
        num_elems = m * n
        r_gx = (num_elems + 255) // 256
        r_ka_bytes = struct.pack("<QQii", w_ptr, c_ptr, num_elems, sk)
    _cache[key] = (C, buf, cfg_ptr, gx, sk, lds, need_reduce, r_cfg_ptr, r_gx, use_aql)
    _gc[key] = (buf_ptr, asz, cfg, W, td_obj, _r_keep)

def _aql_dispatch(gx, gy, gz, bx, by, bz, shared, kernarg_bytes):
    """Submit one AQL packet and wait."""
    h = _hsa
    # Copy kernarg to device
    ctypes.memmove(_aql_ka_ptr.value, kernarg_bytes, len(kernarg_bytes))
    # Reset signal
    h.hsa_signal_store_screlease(_aql_signal, ctypes.c_int64(1))
    # Get write index
    wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
    slot = wi % _aql_qsize
    addr = _aql_qbase + slot * 64
    # Build packet
    p = _pkt()
    p.header = _HSA_PKT_INVALID
    p.setup = 2  # 3D
    p.wg_x = bx; p.wg_y = by; p.wg_z = bz
    p.grid_x = gx * bx; p.grid_y = gy * by; p.grid_z = gz * bz
    p.private_seg = _aql_priv
    p.group_seg = shared
    p.kernel_object = _aql_ko
    p.kernarg_address = _aql_ka_ptr.value
    p.completion_signal = _aql_signal
    # Write body first (with invalid header)
    ctypes.memmove(addr, ctypes.byref(p), 64)
    # Atomically publish header
    ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_DISPATCH)
    # Advance write index
    h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
    # Ring doorbell
    h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
    # Wait
    h.hsa_signal_wait_scacquire(_aql_signal, ctypes.c_uint32(_HSA_COND_LT),
                                ctypes.c_int64(1), ctypes.c_uint64(0xFFFFFFFFFFFFFFFF),
                                ctypes.c_uint32(_HSA_WAIT_ACTIVE))

def _aql_dispatch_with_signal(ko, priv, gx, gy, gz, bx, by, bz, shared, ka_ptr, ka_bytes):
    """Submit AQL packet WITH completion signal and wait."""
    h = _hsa
    ctypes.memmove(ka_ptr.value, ka_bytes, len(ka_bytes))
    h.hsa_signal_store_screlease(_aql_signal, ctypes.c_int64(1))
    wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
    slot = wi % _aql_qsize
    addr = _aql_qbase + slot * 64
    p = _pkt()
    p.header = _HSA_PKT_INVALID; p.setup = 2
    p.wg_x = bx; p.wg_y = by; p.wg_z = bz
    p.grid_x = gx * bx; p.grid_y = gy * by; p.grid_z = gz * bz
    p.private_seg = priv; p.group_seg = shared
    p.kernel_object = ko; p.kernarg_address = ka_ptr.value
    p.completion_signal = _aql_signal
    ctypes.memmove(addr, ctypes.byref(p), 64)
    ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_DISPATCH)
    h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
    h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
    h.hsa_signal_wait_scacquire(_aql_signal, ctypes.c_uint32(_HSA_COND_LT),
                                ctypes.c_int64(1), ctypes.c_uint64(0xFFFFFFFFFFFFFFFF),
                                ctypes.c_uint32(_HSA_WAIT_ACTIVE))

def _aql_dispatch_no_wait(ko, priv, gx, gy, gz, bx, by, bz, shared, ka_ptr, ka_bytes):
    """Submit AQL packet without waiting (fire-and-forget). Returns write_idx."""
    h = _hsa
    ctypes.memmove(ka_ptr.value, ka_bytes, len(ka_bytes))
    wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
    slot = wi % _aql_qsize
    addr = _aql_qbase + slot * 64
    p = _pkt()
    p.header = _HSA_PKT_INVALID
    p.setup = 2
    p.wg_x = bx; p.wg_y = by; p.wg_z = bz
    p.grid_x = gx * bx; p.grid_y = gy * by; p.grid_z = gz * bz
    p.private_seg = priv; p.group_seg = shared
    p.kernel_object = ko; p.kernarg_address = ka_ptr.value
    p.completion_signal = _sig(0)  # no signal
    ctypes.memmove(addr, ctypes.byref(p), 64)
    ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_DISPATCH)
    h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
    h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
    return wi

def _aql_barrier_wait():
    """Submit barrier packet and wait for all prior work to complete."""
    h = _hsa
    h.hsa_signal_store_screlease(_aql_signal, ctypes.c_int64(1))
    wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
    slot = wi % _aql_qsize
    addr = _aql_qbase + slot * 64
    ctypes.memset(addr, 0, 64)
    # Write completion signal at offset 48
    _sig.from_address(addr + 48).handle = _aql_signal.handle
    ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_BARRIER_AND)
    h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
    h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
    h.hsa_signal_wait_scacquire(_aql_signal, ctypes.c_uint32(_HSA_COND_LT),
                                ctypes.c_int64(1), ctypes.c_uint64(0xFFFFFFFFFFFFFFFF),
                                ctypes.c_uint32(_HSA_WAIT_ACTIVE))

def custom_kernel(data: input_t) -> output_t:
    A = data[0]; m = A.shape[0]; k = A.shape[1]; n = data[2].shape[0]
    key = (m, n, k)
    if key not in _cache:
        _setup(key, data)
    t = _cache[key]
    C, buf, cfg_ptr, gx, sk, lds = t[0], t[1], t[2], t[3], t[4], t[5]
    need_reduce, r_cfg_ptr, r_gx, use_aql = t[6], t[7], t[8], t[9]

    struct.pack_into("<QQQ", buf, 0, A.data_ptr(), data[2].data_ptr(), data[4].data_ptr())

    if use_aql:
        try:
            if not need_reduce:
                _aql_dispatch(gx, 1, sk, 512, 1, 1, lds, bytes(buf))
            return C
        except Exception as e:
            print(f"[AQL] dispatch error: {e}", file=sys.stderr)

    if _fn is not None:
        _hip.hipModuleLaunchKernel(_fn, gx, 1, sk, 512, 1, 1, lds, None, None, cfg_ptr)
        if need_reduce:
            _hip.hipModuleLaunchKernel(_fn_reduce, r_gx, 1, 1, 256, 1, 1, 0, None, None, r_cfg_ptr)
        return C

    A_d, B, B_q, B_shuffle, B_scale_sh = data
    import aiter; from aiter import dtypes
    from aiter.ops.triton.quant import dynamic_mxfp4_quant
    from aiter.utility.fp4_utils import e8m0_shuffle
    Aq, As = dynamic_mxfp4_quant(A_d); Ash = e8m0_shuffle(As)
    return aiter.gemm_a4w4(Aq.view(dtypes.fp4x2), B_shuffle,
                           Ash.view(dtypes.fp8_e8m0), B_scale_sh,
                           dtype=dtypes.bf16, bpreshuffle=True)
scrolls · 707 lines total

Source code from GPU Mode and the KernelBot dataset · June 9 Researcher Reciprocity License v1.0

Changes from previous submission

Against this author's previous submission submission 728109.

  ⋯ 347 unchanged lines
  _aql_qbase = 0
  _aql_doorbell = None
  _aql_signal = None
- _aql_ko = 0       # kernel_object
+ _aql_ko = 0       # kernel_object (main kernel)
  _aql_priv = 0     # private_segment_size
- _aql_ka_ptr = None # kernarg device pointer
+ _aql_rko = 0      # kernel_object (reduce kernel)
+ _aql_rpriv = 0
+ _aql_ka_ptr = None # kernarg device pointer (main, 256 bytes)
+ _aql_rka_ptr = None # kernarg device pointer (reduce, 256 bytes)
  _aql_log = []
   
  def _init_aql():
      global _hsa, _aql_ready, _aql_queue, _aql_qsize, _aql_qbase, _aql_doorbell
-     global _aql_signal, _aql_ko, _aql_priv, _aql_ka_ptr
+     global _aql_signal, _aql_ko, _aql_priv, _aql_ka_ptr, _aql_rko, _aql_rpriv, _aql_rka_ptr
   
      if _CO_BYTES is None:
          _aql_log.append("No .co bytes")
  ⋯ 112 unchanged lines
      _aql_priv = priv.value
      _aql_log.append(f"Kernel: obj=0x{_aql_ko:x} priv={_aql_priv}")
   
+     # Load reduce kernel symbol
+     rsym = ctypes.c_uint64()
+     err = h.hsa_executable_get_symbol_by_name(ctypes.c_uint64(exe.value),
+             b"reduce_bf16_kernel.kd", ctypes.byref(gpu), ctypes.byref(rsym))
+     if err != 0:
+         err = h.hsa_executable_get_symbol_by_name(ctypes.c_uint64(exe.value),
+                 b"reduce_bf16_kernel", ctypes.byref(gpu), ctypes.byref(rsym))
+     if err == 0:
+         rko = ctypes.c_uint64()
+         h.hsa_executable_symbol_get_info(ctypes.c_uint64(rsym.value), _HSA_SYM_KERNEL_OBJ, ctypes.byref(rko))
+         rpriv = ctypes.c_uint32()
+         h.hsa_executable_symbol_get_info(ctypes.c_uint64(rsym.value), _HSA_SYM_PRIV_SEG, ctypes.byref(rpriv))
+         _aql_rko = rko.value
+         _aql_rpriv = rpriv.value
+         _aql_log.append(f"Reduce: obj=0x{_aql_rko:x} priv={_aql_rpriv}")
+         # Alloc separate kernarg for reduce
+         rka = ctypes.c_void_p()
+         err2 = h.hsa_memory_allocate(ctypes.c_uint64(ka_region), ctypes.c_size_t(256), ctypes.byref(rka))
+         if err2 == 0:
+             _aql_rka_ptr = rka
+     else:
+         _aql_log.append(f"Reduce symbol not found: {err}")
+  
      _aql_ready = True
      _aql_log.append("AQL READY")
   
  ⋯ 11 unchanged lines
  _gc = {}
   
  def _get_split_k(m, n, k):
-     k_steps = k // 128
-     ntm = (m + 15) // 16
-     ntn = (n + 127) // 128
-     tiles = ntm * ntn
-     if k_steps < 8: return 1
-     if tiles >= 64: return 1
-     for sk in [2, 4, 7, 8]:
-         if k_steps % sk != 0: continue
-         if tiles * sk >= 128: return sk
+     # With AQL dispatch (~1.5 µs), sk=1 is faster than sk=8 + reduce (~7 µs HIP overhead)
+     # Even for case 2 (17 tiles, 6.6% CU util), the saved dispatch > lost parallelism
      return 1
   
  def _setup(key, data):
  ⋯ 40 unchanged lines
          r_cfg_ptr = ctypes.cast(r_cfg, ctypes.POINTER(ctypes.c_void_p))
          _r_keep = [r_buf, r_bp, r_asz, r_cfg]
   
-     use_aql = _aql_ready and not need_reduce
+     use_aql = _aql_ready and not need_reduce  # AQL only for non-split-K
+     # Pre-pack reduce kernarg bytes
+     r_ka_bytes = None
+     if need_reduce:
+         num_elems = m * n
+         r_gx = (num_elems + 255) // 256
+         r_ka_bytes = struct.pack("<QQii", w_ptr, c_ptr, num_elems, sk)
      _cache[key] = (C, buf, cfg_ptr, gx, sk, lds, need_reduce, r_cfg_ptr, r_gx, use_aql)
      _gc[key] = (buf_ptr, asz, cfg, W, td_obj, _r_keep)
   
  ⋯ 32 unchanged lines
                                  ctypes.c_int64(1), ctypes.c_uint64(0xFFFFFFFFFFFFFFFF),
                                  ctypes.c_uint32(_HSA_WAIT_ACTIVE))
   
+ def _aql_dispatch_with_signal(ko, priv, gx, gy, gz, bx, by, bz, shared, ka_ptr, ka_bytes):
+     """Submit AQL packet WITH completion signal and wait."""
+     h = _hsa
+     ctypes.memmove(ka_ptr.value, ka_bytes, len(ka_bytes))
+     h.hsa_signal_store_screlease(_aql_signal, ctypes.c_int64(1))
+     wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
+     slot = wi % _aql_qsize
+     addr = _aql_qbase + slot * 64
+     p = _pkt()
+     p.header = _HSA_PKT_INVALID; p.setup = 2
+     p.wg_x = bx; p.wg_y = by; p.wg_z = bz
+     p.grid_x = gx * bx; p.grid_y = gy * by; p.grid_z = gz * bz
+     p.private_seg = priv; p.group_seg = shared
+     p.kernel_object = ko; p.kernarg_address = ka_ptr.value
+     p.completion_signal = _aql_signal
+     ctypes.memmove(addr, ctypes.byref(p), 64)
+     ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_DISPATCH)
+     h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
+     h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
+     h.hsa_signal_wait_scacquire(_aql_signal, ctypes.c_uint32(_HSA_COND_LT),
+                                 ctypes.c_int64(1), ctypes.c_uint64(0xFFFFFFFFFFFFFFFF),
+                                 ctypes.c_uint32(_HSA_WAIT_ACTIVE))
+  
+ def _aql_dispatch_no_wait(ko, priv, gx, gy, gz, bx, by, bz, shared, ka_ptr, ka_bytes):
+     """Submit AQL packet without waiting (fire-and-forget). Returns write_idx."""
+     h = _hsa
+     ctypes.memmove(ka_ptr.value, ka_bytes, len(ka_bytes))
+     wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
+     slot = wi % _aql_qsize
+     addr = _aql_qbase + slot * 64
+     p = _pkt()
+     p.header = _HSA_PKT_INVALID
+     p.setup = 2
+     p.wg_x = bx; p.wg_y = by; p.wg_z = bz
+     p.grid_x = gx * bx; p.grid_y = gy * by; p.grid_z = gz * bz
+     p.private_seg = priv; p.group_seg = shared
+     p.kernel_object = ko; p.kernarg_address = ka_ptr.value
+     p.completion_signal = _sig(0)  # no signal
+     ctypes.memmove(addr, ctypes.byref(p), 64)
+     ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_DISPATCH)
+     h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
+     h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
+     return wi
+  
+ def _aql_barrier_wait():
+     """Submit barrier packet and wait for all prior work to complete."""
+     h = _hsa
+     h.hsa_signal_store_screlease(_aql_signal, ctypes.c_int64(1))
+     wi = h.hsa_queue_load_write_index_relaxed(_aql_queue)
+     slot = wi % _aql_qsize
+     addr = _aql_qbase + slot * 64
+     ctypes.memset(addr, 0, 64)
+     # Write completion signal at offset 48
+     _sig.from_address(addr + 48).handle = _aql_signal.handle
+     ctypes.c_uint16.from_address(addr).value = _mkhdr(_HSA_PKT_BARRIER_AND)
+     h.hsa_queue_store_write_index_relaxed(_aql_queue, ctypes.c_uint64(wi + 1))
+     h.hsa_signal_store_screlease(_aql_doorbell, ctypes.c_int64(wi))
+     h.hsa_signal_wait_scacquire(_aql_signal, ctypes.c_uint32(_HSA_COND_LT),
+                                 ctypes.c_int64(1), ctypes.c_uint64(0xFFFFFFFFFFFFFFFF),
+                                 ctypes.c_uint32(_HSA_WAIT_ACTIVE))
+  
  def custom_kernel(data: input_t) -> output_t:
      A = data[0]; m = A.shape[0]; k = A.shape[1]; n = data[2].shape[0]
      key = (m, n, k)
      if key not in _cache:
          _setup(key, data)
-     C, buf, cfg_ptr, gx, sk, lds, need_reduce, r_cfg_ptr, r_gx, use_aql = _cache[key]
+     t = _cache[key]
+     C, buf, cfg_ptr, gx, sk, lds = t[0], t[1], t[2], t[3], t[4], t[5]
+     need_reduce, r_cfg_ptr, r_gx, use_aql = t[6], t[7], t[8], t[9]
   
      struct.pack_into("<QQQ", buf, 0, A.data_ptr(), data[2].data_ptr(), data[4].data_ptr())
   
      if use_aql:
          try:
-             _aql_dispatch(gx, 1, sk, 512, 1, 1, lds, bytes(buf))
+             if not need_reduce:
+                 _aql_dispatch(gx, 1, sk, 512, 1, 1, lds, bytes(buf))
              return C
          except Exception as e:
              print(f"[AQL] dispatch error: {e}", file=sys.stderr)
-             # Fall through to HIP
   
      if _fn is not None:
          _hip.hipModuleLaunchKernel(_fn, gx, 1, sk, 512, 1, 1, lds, None, None, cfg_ptr)
  ⋯ 1 unchanged lines
              _hip.hipModuleLaunchKernel(_fn_reduce, r_gx, 1, 1, 256, 1, 1, 0, None, None, r_cfg_ptr)
          return C
   
-     # Fallback
      A_d, B, B_q, B_shuffle, B_scale_sh = data
      import aiter; from aiter import dtypes
      from aiter.ops.triton.quant import dynamic_mxfp4_quant
scrolls · 181 diff lines total

Best evidence level for this revision: reported

JSON