submission 728109
bhagawan-yantrion · python · License unknown
Use it
Vendorable · source mirrored · license unknownView source →
No package. Vendor the mirrored source: 620 lines, June 9 Researcher Reciprocity License v1.0.
submission_aql.py
curl "https://kernelindex.com/api/v1/implementations/kernelbot-amd-mxfp4-mm-728109?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
Reported · How evidence levels are derived →
Source and license
sourceavailable
revision digestsha256:4a2e47771f0036c572eb31c6eb4efa4bae46d0a5e416614162578c03ceddb1a7
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-k
def _get_split_k(m, n, k):Kernel source
submission_aql.py620 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
_aql_priv = 0 # private_segment_size
_aql_ka_ptr = None # kernarg device pointer
_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
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}")
_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):
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
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
_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 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]
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))
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)
if need_reduce:
_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
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 · 620 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 727321.
#!/usr/bin/env python3- """v15: self-cleaning split-K with lightweight fences. Single launch always."""- import os, ctypes, base64, tempfile, struct+ """HSA AQL direct dispatch — single-file submission for testing.+ Falls back to hipModuleLaunchKernel if AQL init fails."""+ import os, ctypes, base64, tempfile, struct, sysos.environ["HIP_FORCE_DEV_KERNARG"] = "1"import torchfrom 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"⋯ 242 unchanged lines"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,⋯ 5 unchanged lines_fn = None_fn_reduce = None- def _init():+ def _init_hip():global _fn, _fn_reduce+ if _CO_BYTES is None:+ returntry:td = tempfile.mkdtemp()p = os.path.join(td, "k.co")with open(p, "wb") as f:- f.write(base64.b64decode(_CO_B64))+ f.write(_CO_BYTES)mod = ctypes.c_void_p()if _hip.hipModuleLoad(ctypes.byref(mod), p.encode()) == 0:fn = ctypes.c_void_p()⋯ 5 unchanged linesexcept:pass- _init()+ _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+ _aql_priv = 0 # private_segment_size+ _aql_ka_ptr = None # kernarg device pointer+ _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++ 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}")++ _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 = {} # prevent GC of ctypes objects+ _gc = {}def _get_split_k(m, n, k):k_steps = k // 128ntm = (m + 15) // 16ntn = (n + 127) // 128- num_tiles = ntm * ntn+ tiles = ntm * ntnif k_steps < 8: return 1- if num_tiles >= 64: return 1+ if tiles >= 64: return 1for sk in [2, 4, 7, 8]:if k_steps % sk != 0: continue- if num_tiles * sk >= 128: return sk+ if tiles * sk >= 128: return skreturn 1def _setup(key, data):- """First-call: pre-compute and cache everything for this (m,n,k) shape."""m, n, k = key- A = data[0]- B_q = data[2]- B_scale_sh = data[4]- sa = A.stride(0)- sbq = B_q.view(torch.uint8).stride(0)+ sa = data[0].stride(0)+ sbq = data[2].view(torch.uint8).stride(0)ksp = ((k // 32 + 7) // 8) * 8sk = _get_split_k(m, n, k)k_chunk = k // sk- ntn = (n + 127) // 128- ntm = (m + 15) // 16- gx = ntm * ntn+ ntn = (n + 127) // 128; ntm = (m + 15) // 16; gx = ntm * ntnlocal_K = k_chunklds = 16 * (local_K // 2) + 16 * (local_K // 32) + 4C = 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++ w_ptr = 0; td_ptr = 0; W = None; td_obj = None; need_reduce = Falseif sk > 1:W = torch.empty(sk * m * n, dtype=torch.float32, device="cuda")- w_ptr = W.data_ptr()- # Partials mode: td_ptr=0 → kernel skips self-cleaning, reduce kernel sums- td_ptr = 0- need_reduce = True- # Pre-pack kernarg: A(0), Bq(8), Bsc(16) are placeholders+ w_ptr = W.data_ptr(); td_ptr = 0; need_reduce = True++ # HIP kernargbuf = (ctypes.c_char * 76)()- _pi = struct.pack_into- _pi("<QQQQ6iQQi", buf, 0,- 0, 0, 0, c_ptr,- m, n, k, sa, sbq, ksp,- w_ptr, td_ptr, k_chunk)+ 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)(⋯ 1 unchanged linesctypes.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))- # Reduce config for split-K+r_cfg_ptr = None; r_gx = 0; _r_keep = []- if need_reduce and _fn_reduce is not None:- num_elems = m * n- r_gx = (num_elems + 255) // 256+ if need_reduce and _fn_reduce:+ num_elems = m * n; r_gx = (num_elems + 255) // 256r_buf = (ctypes.c_char * 24)()- _pi("<QQii", r_buf, 0, w_ptr, c_ptr, num_elems, sk)+ 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)(⋯ 2 unchanged linesctypes.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]- _cache[key] = (C, buf, cfg_ptr, gx, sk, lds, need_reduce, r_cfg_ptr, r_gx)++ use_aql = _aql_ready and not need_reduce+ _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 custom_kernel(data, _c=_cache, _pi=struct.pack_into, _launch=_hip.hipModuleLaunchKernel):+ 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 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]++ 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))+ return C+ except Exception as e:+ print(f"[AQL] dispatch error: {e}", file=sys.stderr)+ # Fall through to HIP+if _fn is not None:- A = data[0]- m = A.shape[0]; k = A.shape[1]- n = data[2].shape[0]- key = (m, n, k)- if key not in _c:- _setup(key, data)- C, buf, cfg_ptr, gx, sk, lds, need_reduce, r_cfg_ptr, r_gx = _c[key]- _pi("<QQQ", buf, 0, A.data_ptr(), data[2].data_ptr(), data[4].data_ptr())- _launch(_fn, gx, 1, sk, 512, 1, 1, lds, None, None, cfg_ptr)+ _hip.hipModuleLaunchKernel(_fn, gx, 1, sk, 512, 1, 1, lds, None, None, cfg_ptr)if need_reduce:- _launch(_fn_reduce, r_gx, 1, 1, 256, 1, 1, 0, None, None, r_cfg_ptr)+ _hip.hipModuleLaunchKernel(_fn_reduce, r_gx, 1, 1, 256, 1, 1, 0, None, None, r_cfg_ptr)return C+# Fallback- A, B, B_q, B_shuffle, B_scale_sh = data+ A_d, B, B_q, B_shuffle, B_scale_sh = dataimport aiter; from aiter import dtypesfrom aiter.ops.triton.quant import dynamic_mxfp4_quantfrom aiter.utility.fp4_utils import e8m0_shuffle- Aq, As = dynamic_mxfp4_quant(A); Ash = e8m0_shuffle(As)+ 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 · 418 diff lines total
Best evidence level for this revision: reported
JSON