submission 631390

  #!POPCORN leaderboard amd-mixed-mla
- import torch
- import triton
- import triton.language as tl
+ import base64
  import os
+ import pathlib
+ import zlib
   
+ import torch
+ from torch.utils.cpp_extension import load_inline
+  
  from task import input_t, output_t
   
+ os.environ.setdefault("HSA_XNACK", "0")
+ os.environ.setdefault("PYTORCH_ROCM_ARCH", "gfx950:xnack-")
+ os.environ.setdefault("TORCH_EXTENSIONS_DIR", str(pathlib.Path(__file__).resolve().parent / ".torch_extensions"))
   
- NUM_HEADS = 16
- KV_LORA_RANK = 512
- QK_ROPE_HEAD_DIM = 64
- QK_DIM = KV_LORA_RANK + QK_ROPE_HEAD_DIM
- V_DIM = 512
- SM_SCALE = 1.0 / (QK_DIM ** 0.5)
- MXFP4_BLOCK = 32
- MXFP4_BLOCKS = QK_DIM // MXFP4_BLOCK
- PACKED_QK_DIM = QK_DIM // 2
- USE_MXFP4 = False # os.getenv("MLA_USE_MXFP4", "0") == "1"
+ HERE = pathlib.Path(__file__).resolve().parent
+ MLA_CPP_B64 = "eNrlPWtz20aS3/UrZpmyC6BoiaRkRStK2kocO3b5nei0dcfVokBwKMEkAQoAKSk+/ffrnvcAAxJSnL3LLcsWycF0T093T093z4PfxUk0W44pOY563d2reLH7+s2nX4uMhvOdq9Ot7/TjaVDEM7obpRndDbPoiv3ZuVosrFp5MY6TwixCnPA/GE16BzZK+WSyqH9y6H6QLZMinlP7YZZGiyye74p3Tpt+3jK7gA9b1rPdaMoLDYxFih2ltwVN8jhN7Obmy4LemgXpooBK4cwsK+4WNCiyMC5yoCUJ5zRfhBEl0bQgJ0RQNNha5nFySSaHUBYEoucB3Z/v6Ue3++bD233r8ah3oJ6OJrM0LHoHg62trbzIllFBXk/fz8KfaJSO6aub8Rkj5zWAfN0i8MqLsIgjEqVJDj1aZAQkuNcPCjL9nH5Yzl/TcIzVoAFEuh7i7aoEsRngXZqFv4TJlAM87/U3gXyefkgXFBv5KZ4DiIFjM+gvFujB/mYIWZvTV2p9u4RzE7rztISuOe2fwkv6a/wbJU1ZC5L4e5gtcgmw3wDg7AqG/jgXxCkMbVTYo6NLWgQ3UBDkQIfnb8L3MYqWizCJ7jTFDGJ3l3x+uwH27+GKnsHY+Kz4LnVx0ADy/Tz8TBorrWzrrYDY6zds4+3D2kCq3G0EYZ7TrPCMbp+caCBfM+7TecPOnEvGKZVrSOL5wxn3SbalyG/Y1qdHtHVeautt07bOm7Ee6hm8h29Pnxr0nhgtf5JyqW/7x1kaTd8TYvQTpXh9BSJpAvpBgwLZCDpdISy5JTm9boLhbQUDgx/H8w3Q78b52xVD8Ut6kzeSkQZ5kc7yRvZVg/x4V1CEKTfcruBtE7Q/6cSbHPqN8IPlAECGIOc2XFrh3TLypvgUMk8KardMua9IN9vfqDK/pMtk/B6mamU1UWjdI5LhA1KkJKFhRvOC0BVNdrbIQ1+ArefAFt6Ed4/D1jew/UazdA0Sx8BD18W07x3S+pjM7ki+XCzSrMhJggp70jv4S8uwg6MwpwS9q5yVcDfoGviHLJscDozS6QqLy6XpsgjQBxuhWt8D4oLOF7OwoMcgCRJ2UCBkdLplSYjQ2zAqgnG88nzLe5KdCckTMsL+dKEbIZkvgbEjCqNtFefxaEbJ6I6MWkJrM1oss4SEoDojToNuLFwCMyMaz7AxIHN44bEiIIy9j2T7Agm0vA0tPyM9X6NTXUJGofNJzk63gmBMV3FEg4CcEXAVx94qjcekvSiyEs62d9b2oRiQbcbFkOTgMVONrwMtrMLZkkrEEiN0iJUzzE4n9VOYhXN0Us97ApaL7QxAa9xZrRxxslgWXDHQM25fB6wX+B187DYohPouh90o76BFndEE3hezuJiuND7QFRvhJAY/P+DFAi863aTNQOsfzHKqsS6wi7SgGUfMZC+q5+mkmIe3QR6FGB6UHrdRpdkj3g8UNeD78PHs5RF0oiCzcJlEV2SEgzLHUQleNZlAMAMGNkrnC5i7MiwOmaQ4Y8PLRaaETLSUX9UzHOV/OUtHwIoAgg/easBbDQIPGu2QV0dH2g/0t1iD0ySYz8JgBYyf3IxBhRBxcNU7YGEP/l/1vHp1YHzLUaW+iydjOoGmX7z74cPPPwVBU0X5joIkSpVfYXEyjidbBr/ZWGMuL9B9UvKC47H0gY3KszChpcqiyKqsB/kUXWzm2J+s8bI1ENeQ6bv0so8gvZ39/f7BX59397uHh4d/Pdjbn1QURukL1G9z9u3YKlTuhRgCUF9Ur4yJfI69zG/i336bUVlIQ1C7y8WSXIU5OSS3EVhzVgafWNk+ya/CMWgfTS7jhIqnVpmAjZYMKTd2OHNyprKPb8a3O7cVzkMbRh38uKt8PaNatCzVeqJq4byim+IV20DLtsBd5dNNmk1pVkLYJmeg9Sp42pYKJCzOhHga7FRxeMQcG85maS+1Meaw94rMnByfkGQ5D6JlXiZqFBYwDuPxLRClW9rVdq1UH1Qhulom0wrIExsEpc5MpH7AprpafFhVaxAHXUsGTKZZIciwqGpbSF2gMHQ1oMazXQIsQ7KpQRPJv24bLOTC1JkQ8X1qBVU2KRZG8d1GabOEo9Spk1ITvA+6SxuqVyhKbYLMuatBT0vhI+jA2cefPh6RcbpEfya6otHUaT2CUrvm1Gi1q9V+EwmmNBtTXKELTL+LMigusaOOqvVUCB0LArBqGR3D1LiUTkYeXAchzOtD0z5cDCr1kzhNDAsAPss1A+HR5MXQkreAFxWnKxv5sM+r23GLUabClXIZRkMC9z1QPhsj9WkWyJHLOPpUPABmWlV2rquV2HRi1wJqhW28MOoDd7ESl9Hw6YWHnjfHgAznrnmW3vAPUTrzybNT6Hzb4JkhdagZMNsM2OAzWKBq5wcuQMCsAOFzBVDHi2UDLGaCWw4okLR19HCspPmh46DmFKKLbU33oDQXaIYONVcuhqrViyF28wlxyRj78YTUS1pP8EOhbh0jXr4gC5inUWTCmb7dB5oug+uh3TUFwHunkl/QMdRI0zO+lk6ulj8GJUylUPZ+WagqFpt+PqfRyxmd5zq5WVMr/0QzntqEqjWklsdxR87iMifaNlo8lX6ZIvo6WNEoEP6QobdKT/Ex09WvVnwsRCqekn8igEZ9b3QLfUNcRbndB1PCuB5AMDIsdfDCqZFiZheObNtwOrelx6rhvltk4eU8BDuUpbOZti5pRliXYkDVHcDbcZm7A7K9HZd7aJCBveSk8Lk51r6S4PKgDhRV+kQh4CrMFbQWBPmpQXYNEEuWg5I8JGPjC4BGXTy2WH/qPWVeAhtk1akaGr0wdONefcJZAEIcoG2UBzAJMCb+C3h+HSz3+lZP2HR022ddMbvr25yoIcki6wsn6wuQ1UdCvpQJKUdHKc4BPGUShXlxLLly6jE6h18uyFPSvX0Fr6XPPn6PHwfrkF7FtUjLpXt9symfnJ6S3sHadrS04K8nvnXYl1nagbb9EtvuHbLXXtNiATEQjji0pGAeiytwAJfgevQImPKbLIbgex7ejejfLMVh7oNoXMjT3TVRx6BJgSc0vrwapVktbBCMlvGsiBNwUsaXEUTp6SoAij2rfyYsrvV6Bll+Bxj5Yw//vhJ/JiE4ViaP7F4YTMWSjqLSgPhXEb7/8k9KeK+//6el/PDPSvl+/09L+d43olyBYE5DJb5gRijPAirwkd4+IDaQuqZLiIfukqjgXoFnNmbTZ8yhm6hoMqEZ86wZRzkn2QbMUh2vzK6uacJGVhE3smISZ3mBnbTEanAHppqXh++7hOf7RncFLfmDejpB6Al8HMVFbguEHB+r5T4Nifxhr1Ec5hCpgo/Nl6iQc77neTbGbdCvfhdfSzbR9vd8Mbcv/YGDqENcpuE5SSRboBZP/HmceP3n+3zm7XZMGp6Rvl9iwa+s9yhVCKrHoHWUvDg/O4KIP885f3hMxrk0jjMaFTNj+Qsn7CsKPeaLExiPAx00S8LZ7I4t54xpjos5/X96DMezXv97GF15CuHFIk0osGoS0xnGG7pLO6VeiwhoGSaFys8GAfY4CPOAPfU8KS/fYA2IB7ipmXitwA1kBkd0oBZMDtdEDGI3EXq6fbY6FrB9W/AWFkUWj5YFDQLPo7cFRjsFBPG4TuAB+/0/zJUVKc0ce+dJ0tq+9lxL/j96uNwBZH1YhNEU9OQmzca2tcCX/bAy4KKVYCXq9GLKliaQAs+A6zDahhjJ9C86Jv/Bvvp/aHMwwnrlJnu+K5wBobNohgcyccFnB9SKUxO5Fbf8vwsHZYjtDgkHpUkUFzJldgtDvYuhFeF7LK5XMb0ZT4JIDK77D5vcHsDsximXTWLgvJQz5wNCawlihdZcjPsNxIB/QY0dCgt8G3ZV+M319JHC8GvQ9x6JHgddfRP3VhZNZ7N4SnNNOpN/9t0ZTNzBEM9k1lqtKJQ8vQOWnOCIYE401wx89yJBOT2eRcEN6A3YpcmEZuBv5Okyi8DwVi3sPJxSZ2XPILZDPHP1pUKFtZPHr83gvV0Bq+X2IAFgJDKg2xtg0SL9KBK5duqTY90122g+kZlpXhiT+usxqabccRCqGusHo0o6a4y7FkZjM1H9+HRgefJALNwwaIrajpSyTiACkbZJsjPMHZPJVv7UaVvMvHkVGbTqaYP1kGZRCT1TpXbdOuhwmFfpZILbK+riM682I+jXaboaT0qVSQbDEJuoGYw2rKIN7Uk4Hmeb4wbmUMLghFGjF1o80x4xc4dUlwgN8znwYBbi8PZaOQs/RxAszMEdf9JtkaMj0spbniRlLXglkGqJfjJzOUZvBFzWJ4gZ00kEWE8TwgVw9KRPuAltVdAgDauWJ0QFHjlShCzFj0nL6/pVENKa03ma3dnYXFZdGfWbEJwobdQvjkcpaO30HSjEGwgTTq1lC4hLtVXyVKWKEUA9yANEHSXFcdcaI/cEt6ysh6gaHRtFdW6C2bDPJF8YU5M5LV1XzJVKZTIXiUzJJEtBroAJdxQtxMJUwawkbiXKoJFlTjFvQNkOwSILk3yR5jw8G9EC63ABkDCKKARoixALE9Um21Hmc7oGpVJOob16xReQZN95lIUT/3zoWIF7L4as3NPd0YsjuFxVxpE/GodjOnHl0R+MvpJ35139gg7OszcfXr358ObsPwfW05w/7e50J+bGEt3LdF0v1QJ7qZuA8WtXjBLD68Fo3XJGTA3E7WgYVQqvgathZTRt9JC0+QqiZca25br9IMNV/QAzR3+wAVsCsS6iKzXAnMA6cIQRlbm7xqMb3bCxYq84fj0d1i8PS8Gz9dOH658UTI1F+kutSTKlWGJKx+ymM8gxbOSO2liYLmgWguXw/GPTWtrZG2ZgzAKkg0BAf7BPJot9sEwkTSgumODUFowO9l2V+zzzYj4aUxYsd8h8HnZIDtwiXNxb9mIu7zdXRuL219fZSDHbea2BxMQ0sOVyL0xXlclSrmg7XK322g0ITpTSyfu1yOIxrax+r9+D0LaJWtsUeDzhbGJu/rfC0IErRy02TZuai4nbKjF+Q/DPb1lepIIDKTJ1TAkaMF3KuMq9q6Ir/q1zkgPwtaS3bYfRslu10Kw683V5hssOqRVX2uRwPQoBrnbLWcq03WRbv9UPt8/dBI3VoVqaYcwoZ9q7ljFOX6hcSWkbtSuxPBFYSvo4qMoere7UWhjHg5WHYtlYqMW2wd5tTXYlq8A2cTCM1uaNxrGkEUfKGLLRdICTf31IaQ5PtLR5/fYXR9qo6msKBvGOBos4uVykyaXa0WO6GMZjlckJpqV5tCPtaLc8Opus1VwbTLuWXJM9Rca4LbO9oAFhh8AiUDyVJyvNV8mFl9w8RTfgq1M9cTFKYARnwSCsDsDBOEnbP0lvMwt1Y/7A2cD9VrWkup4lRwagm4AISuODx6hy38h0OJVG0l+7Q4P3S+Czk9AqceXuewPE1xahCh8ncXi9jkS7o9X11ck8ZOtZ4E3APzyvwA8oeFv10kNkndrn15uer0PQfVixo8Mir12WXsdse81CZclvBQC+NcYRZjw6kb3G0MWPWdH9pqGVLSrMI0Mg1dY7ic1zO0Sf04GP1+ZhHvOlmAjKdjvx+NeO1cJDdhixzUNsnxFuMNra0AruOp7MgltwyGVR78DggZnDaY4CV74tFFW9yZdzKwD9N1MZnIYX/YnnWaXPBI/5yVE84OM79QWZt31iYdysIRrOlhgUloU+aAZki9mQshUL3wRzpTM8M4H2RnTUGT9jrYzKFW2erwAf4pSpC/mbyTpEBmxjrRhcI0ekZ6mW7XpK1AKRUOwqmkEppcIcGVapnE5hT9THttWFbcnFtmr739VMVvtfO8MwWaEJCxDJgh+BVVtXOvrgZrldsIThERlm3Z2Q/DfJevytz9/2dsKLcu0Rrz3itUe89ojXHl04Ihh7Ey0JR+7tOkg8Bg9AZn4TLjwgeiS2W3VIkS2pc7VBdZEhRdw1wadigK7YK1VEVrCnZW7o3la5IQA0wx7BEHEdTQ1DegeCIWhARvinnimylwvWx95BhRmywrWs4GDCwmRB32JB38ECdvRAsmpPMs7NABH52Rz4ugADd+8OBCeHbFQs1qXx5Nh2JJF/v81Y14jTenBy2f4Rc91rcnjqYRP7YJHFOO9eYLUj+bUHX8Gq7rv3QBiuKLBLj3RPuvF7eof6At33DnE+AfSWi1EJVFeFO1LVmb6Vufyx2pDNI1C7Pp1nHgIGjcftZmZOqLJb6dunAGuTY+xek97Bxlxa+RYPvMbjUckzDE+bJdCsxBcesGeHbi0u7pY3ijz6hNf6/F3d0amNGTy2SmBSLDJRzRJ3nrcSSbB9//Enxfz/E6m/+q0/RtJvJbJ1+77tWpzXQheZ0LW67JxEX9JVvqkR1VJPUSflLY24YE3zPGDXs3l7YFJKNr0mgbgqKhkNO5VYmQXlqkWRHeLKRbDqx3t9z7Poa/u8s44ldNMM0WTcqs9Ergr3DFSyOefuxKCxyGdKx8oLlq2rTvStCq/7gC2hK53IW8k83rlK463qt8zYGbyVO4FXSt6dr83dYd5uZabtztdm7coscKXsgBmrNbm5+636b0rP0o3ZuHS4Esn/zf77ecm3cGySS79FVsyRkXMwC6b2JpDcCXDGGIJYK+6q1uw2K/KdW0F/L4s7BqG1mzL0csDKcZnBQBYfmxcePDMXltVq84lRauqusiDpgozS8V3LCsGttXjnoi3z1OUuQ9/cPCAxhzgsYci01M0lm7Gi48+QujolfQ8cmfJIfmnwmiddkxXzI9mNZZgv2DXjcpmQ+12mqbGhXG+7HrT39puH8OcNQ/jqeVrj0gPlWW679wK3ITrYJl8ceWh9Q4Rxjta86AHP0ZY2nloj/vFDESlCXVCasjl3Zp4sOjY6WDnnZF4wMRQQMjXE8kfAqPSSJ8Jye01CWAL3jVi/0PEyotVLsaxboaqXTzkez+SqLMTKbVK9yEpfiGV/V7e2oK3CK3HBy2i9aBHjCijzcidRHRiMhOPlTt7a7uhLndZcZlPa8MzDQH0jUPUyJrw5rhHsXRW2YJf6OMH4xv1Km4ZPleo7Zmsuoaq5qETjENcQ5GtRlG5FLaEwbpaSNzGWarwPb39FSb1diYsx9c6soalOHa09HUtpOvzmNC7tC3WnCm+M79Ab5dUrUcQ1Z9wYDw0yLgxjX7DRpjPM5pUoDB4HWu0ClDKLnHKuBYByYBQcK7U2CrdP3G2Wo7ecCqSe69KYbY2RRWVKnNtKK227ozvDc+Xiu8H5oWjSnU/ZNJGp0EurBcbUsvwUpw7xefcEQ9evzcgTCwPj9CbRhRxRzUJSZR0RwVwLQn9GEYrlEhhiE68qOXBsBI/+t2RoL+coqbEFnarUnELTqsCxaLF33bBavgYAa7EG4M8meGnIFJhY4dugAiA8wYmqG1B/TKw819BI33+jGMd9tYJf9gJVBqzk2LwOmpVsi/MhULBzK+pWDW0YRUN8wneSwhjtkPLf+w3S69bIbnvb4HWtwPC2Voan4hB6j5UgKy95nLW5pxs89M5OflQkPfg9Z8y5aGrOluMcyE+cxZbk6nIRKKX4AkVqTt1DyTtAg24v70rzNMT9ow6C1vXrIX0y/YyhJeam8uR9dkcRpmcHvvCpx9m3Ya8F88+/G9MJ3lT5+s2n4MXrly/eetF87JPHv/7BcI9T8pV8q9c/FOFX8eJlluFZbUKzDJwvsHpI7+DxKM8+/vLitei6wHmC7fy6ZAc4OqQFrMHW0uyItPCWnMXPtGBU4J7O5JJDgVMgkN6Tmyvc/o/nc7ag9qtlwk7d48kmTKFd5WGUBlMIOihoRBLJS4Z5+TgsQn7ekB0c4DGHfUEzvI2PjthvdpD50rrlG5p7n46XMwqNzdkHXOVfzmbqkmBd0aALo5wJfG1UWYRB7vpAGDuid7kMs/GxJvSUsBJvLu9LwMEjKVRYzJGjFVJ16h2w7yfgj/eUQ2Ieil1lsMioTFgzO4cMbZ96Nketg/0O7CBV2UtoQHCkQ2RTreC/9g6aXPf7pr9XE+S8XPWf194F3HoIeYYMDApr41WF+37LSqzjc7bkqSV6pJRBXGgNEuT3I/CfkDk6OqNJnmanrF6VAZ5Vi1x3Hn4NO7FRTFffAkdQk0jd/BKZWzkfPwaF4ybsx6BhwpA/zVMWB1N2OYBMo3a9M47nno/DbA9/AOJ6h13G3PP5GqpR1GdFz78/AF26Vhe+DyEq7h10oPhC5kFN9NNVCT8UmA2YRWYD4NmoFoq0gNkReNtb0wzn206ynNMZb67H0IgtUAwZDP8ijBOeSJvdsaM1dEbnNClcSKHfABuKqz8YTsHV6Y+v+E8PWayQv0dUw4daZAwX/xWkSWL3HYYM4eVre70eNe47KXECUfNHa/HGeRAtx6HnuxCkCXnxHz/9gL+jRfjl+C5cRh4bj6iKbxIHlLb86l3Fo7wuISV0ses70lhCTepApZopWJNMaPGUqwzzvtgZZk3lCSezDKRafMIo5r+FoJup+VEEhd/Vc3FR8onuza5KUFpsdV1PjQzm5c6WRdV1UnpC9hkmHJZKchBzOUSHv/kwK+IF4E4n7IJuvPIqb1mHHrk5yvndOie22f3In3n+DtceNEX8gw9FTJfLw80fVFFP2Daux2IWA8REbF0brTDT+aK4875iDhCH/fNe/75jdc9EYd0A7UIhZwsHrolNjb6yeQOeNSjYr9zhrrHi6AjHM7t5/8Uyy8DwqV/BA27l4oOQ4KYfJjDP4ZYcLUy542CV3hUwvlNfe3LYZsPTWRu6qXdS8d+t2Kq5AY6Pc3mzfGcTeaag3W2bclQ1jpnlPPUq9UBIayuZW8r4U2vON6oqA1xBZy2irl1nEG4bF1WQjr58/WO69XjmOtcpcr+heKtn/LlDz3fOsM87V2EesJ9fgVnxb7zs2akmBtxZK0DBqEMhsdaSET16tHizw0ldoKZA5VWDpR14QiJNcZT2iPIwEIngIh2yFNVT/uW+XM+Svqha0QgDCn/xAP3HCCSGl6613rHfOMH9RO/f/UC4J084peQmLq44N1tYlwHSZDxbG6FwhG8ZBk/wsrGb238OdrKH/5pC1K8k6Tu3HogS5lduITtEi6GjolPfvp+xZgG/CXe0ojwwDsApocES2kM6zV/s9tLew+EMllna5+Taffn296/2+lvNMt29DF63RA7I5OXPP7/zvCZx+fGpX9s3iGD2PFDBDTWaqFw9im7tE8HBusdiQbn0qxUifLc8kTAZG/4EKKmho5s2njj4WptQWKskjFUNFdXfjAk0c32t7tqnG3jLlbFkNl3Zksaqer/1P5XKb4A="
+ GPU_HSACO_B64 = "eNrtfX98HNV17927s6Od3dFqtJJWa1mWV+uVkIUtyz9jE0gkEIlJTWIopDbYEbIkeyXLkiLJxuRjdlc7+m0jOwoBQ11M4rSBRBhT2iZpeJIghJA0SSWRBmhoAk372gQHaNL2lfcSz/ueO3el1UYC8vLI++NFn8/xV3Pu9577Y845997Z9Sh+7Y4PcIejRmH2j5O9zBxs/qdG4m3NNn5EJd1WpuHfXJbDcMmUNF4m/pdjIbql3iHrLfUz5V2IzJiv56JfdKkvzsDrMzCtHvWV7Zb6OxZiZ6rhrMXrdcp6ncMLkaWNJ70ef4ftpTA1Lzf+c0+T8hv0M9XuDainst/8R5GyldvXmfiSuhCVtOafgS4ErL2+7oM7b2bMPN9wqCna3VB1sLmrvbmte+yHzP/FqoYDnV31jR2H23vYn1Y1dB3ofsD8fFXH/v3dzaTobvlEc+iLVUca2g431x9saW96aN8d9eJqjhWyWUo66y+iLU1Nze31+9o6Gg/a1uuPztWIvMMad8zVqHiHNT4xV2ONXYOn13hc1jjQ1XG4s56K0zpV9c4qzPep+p1VmO/ShkUq/Lms0NV8qAGXzV1pPdr0jvjzHdryjvjz/dlu893p/L9KDaCtY19DW71NTOvTzndcZ75fu95xnfm+3bbIWC7MTW5LU31Ty6Hur1TJeW4+cKgZd39/y9HmJjHr3078lfBy+PNcaUNby4F291/+mp7433GEHq5qa2g/cLjhQPNDH+lsbr9mR+iax+ZU9Ueau7pbOtrHOPty1aGGo/X72xp66m/v6Do4f6e/oyBe2hsONX/lYHv9obaG+u7Otpaeg0cw902HG5vro+u3fLWqs6vlSENP8yJdZl+s6p4Lxc2P2RfdnS1tbTI6P1/VfcehfR1tTyxpvupg05erDre37O/oOiQ6l+aHjserDnc3d9c33YEutjTWd/c0NB58CtM71+bWx+yL9DYfrbq94Ujz/q4OOU81i6eMT2amjO2Lpoyl79d3Ga8xfrNbtv23v2Fc3rAX6+tv2bhFTuuRjRvq998OB2tu7GgS01q/v3OrkCPrr9uwcfvB69sa6kThB25vuqmroaWne/v6Ldce2bB5YdHOhq6GQ1T00fXv+L5H3uK+v7jvXe7lb+c+3/6Pd+I/j8glqAe3s7nnxRAuDzS2rwWstUvWrj2w/+i2zdVXHG1HC2tTS1bqdjr4W6+XtN76nVlz+41dcr1fDsl1qqzGu3D/szXFX5+mwM9pyU9k8P8X6Zlnbn/hSGs3JUWQhKu6NMQMB10PzboPN/2PT1z+5cs/vfNu3lH8YmD9rCr3H6n6Lja/L3G+xV6ALRn87K3SAnuXXYe9+67J6uujLZ31jYeR/Js3bV3fvHnrvn37N2/ZXL1+w9L+kJAYkhM9ZSzNzYYMOB5lfY75/Z5RYzvc9pS9sqXrXwUpcaJ+2n7xt/5xeaZEu4oqul7j5lPk2lPs9J2todMTa51Dk/GpE5OlbOhOi/MYmJab/WKGmf2sNTQ0cTcbmqxRlAse9u1pxeDs/fi9lBePKeyb4jpVhzn1mBI6MdEOfjHj0T5DiTjYr2YcWjazFKWc/dKyhtx9Ew6XElWkLZU9M82dQGVg4ElH30S6LTU0OvFHsDVg8FAvd0enGK+w3J7yS5aVVcz+ZMLJHgDfJ/hJdnYi4bh/gn1ciVncL3QrVB4tVvSoB31Q0Qdd9OUOdjZYFD6bHwh7mHvW49oNvXtWB7LEzpd0FtcJPSzu1Wg87FvTKz2crUyu7M1y8pia7Qs4nEUBS1Wj5BcYl8CVPD/28ZwTkyrbPq0v40xf5g/A4CSOMDMJvox1yjJvIWfewuJAbzB4k1eWeR0bZ71Z64tV9sh0IkBz0TcYU/oGlKL7B41Y37AaGxjIyvMHslf4Aop6ajBWdGn4WWt4Ejj4dL4ShavNZDmvppif0bKvRox/Z1rxcsoDMy4X9Fkq+swDitE36MjiAeSMWZ7tUJRY30jMd2nQExsYxu8DqjIyGNNxHRqYCNA91JUox5xxzF3CMTJxVlHDCuaKy7lSWNw5petRYZt9bbpU52I+KMkonpFBGrNoG23GPJeGX6c+ey4NxlSI+9IgZydmn0T20thds08q6K9j06xW+yvuYZdmeB4mncVe8DgYd5I+XtZHPnNK5YzaI3/lLgXtsplTyKYOFw94UU+R9byoR/wh8NO5Q+Dqkkd1iKtLbl8Gtw/c7AxutuQmwZ30u8tTfJr7JPg+8J1pfJ/gf2d60j1v26kUzyT8fpZUSyIpXg54qXLERyCBOGgNjUycwX2YcioXdBkvk86FsUc85lQRd6MTB+biTp2PO1Wdj7ssdUHcOdPiLmVHDZ2auGEu5kJRKxQW8Uaxl+DuWCruOLgUbxbibuf4+adynZEY+zjixemJFevl0RU+Hs1HH/LQh0Aq7ioqw2fXVIXz4UP58KEAMCB9KZ/F8wkDLF7wVCQSpbhLROBPPp8dZ/AzQooJ8qv8q//8uXwX30UxZlVy3PPE1VZFuUCV1U2vrODIn8BKyr2JyTzcn/zLimitnClwhilned77hS+/3yovF+io5bG82nLkjgrP5z78B7lW+RqBBeA7VodZPjBvbRFTatCuxoMUg1+XMUj311jFKf5mcymuvJyL2DIuDepzsXVikGJNRz4tp/ukujHGX80oIrZOYB5HxTyedXvCKuZFkfOisrhrCrlGpfiCz6n6iUHL7RYxRvPgcnkUy6OLcSsYrw/u78B4XRT72X5meTxifK5aT0zMlUcX40qVu5y48fql4VdoLLodlxSjMvacSb87gvbK7U1G7AUK0/iTJyYR06k4c5Lfpq71jOvsjGtfxnVOxrWCdBBrvTQw5MCaA18vSq05wu/LxxRj7E7mpNg4M/GH4JS/nb8bp+9UpK8rrbavU3019MDE+x3k50boScNeW3rdnpDwa/b5iV5PfrTfPXZnghfBr91o77MT/8jsdVK7Vo+VWxW95I/lVmUvzU25taY3AXw6P1BewR6eUJ2+WCX74sQa9shEkgUjCcfncI/Pi3tM9so9PFruLoqu9pREK9zh6OWe8miluzK61lMVXePeEPVhPNkYjyHH5QfmAimmcmRMFQCDMsaKgMuAxcBC4NlQOHy2uCQsYi5SHj5bvUHE3tmtV4TPbtoS9sHHfPAxA2gA/UA/MDM2g8AgsAhYBCwGFkvf9LG4j9Bg8RxCP4sbMpZzZSz7CYMsnkdYJGO8GDH+dCAQvWzost5E4O47WyOfG4SfRmnBEr4NvAzzfgT+7UOc+1JxXmLHeQ/0BvQG9FbxGqHrgs4PnR+6quKhY+k5wkBcVJUgl4Sr5uKkKozrUPVcvqgK4bp8g7imvFGFVGJFNolrD11TLqrcIq51ukZesSq2imsfXVdwkVtysA771hazHPyeu85PjwpnjPVFWBmQPzaGMc9sxr+5klE+KnjPBpFX8rddMZ+PiottLCmxMRSyMRy2MRJZkLesioq5/OWrLY7l1FbGjNpQLLd2TcxfG0FOq4rl11Ygr5XYea04bGNluY2hahvXbLAxssnGqi02VmwVSONxrMTYVhWLsRRcNj82GksgbWyUMwvTxuareew5XzYPGkAD6Af6gfnAfCBzlLyg+3Kw0AONXKa4gP48pmYB8wuYFQxGaTMci1waXk95KoIcFb40aObz8gS/H3vYgYm6tNhMMh/ibWiC9n3ZuSsUikWLjYjYK3cHo+WeQJT2gjrihOLMS/FCe0LES2ovmIqP9D0h/D37Sc6jCQ6fdQ8JnxVrE3yWkHxOh8/pKX8tCtjrUrDY9qPg0DHyPU+RvSbpK4KM9oDZhcjTgYB9P4NBgd7aQEyvDcasQJF9H4LFAnG+mPEu9zOqq2P+dDl/Hh25Lz9fxA5yuL02YZ/Va6hRS3WLNdyp3X8nfHRWwRrVSutTyaXBorn16f7BWPEl2lsOrJ6bSzVWzufXqSTXMa8DyGNDc2uVO22tctNahfkRe0GEamqdUtSBwfm1ymfHmGfomFincjyLrFE+e41CmSvL48Fp+AXVZe8BKMgwruGXaXzBS4MJplxQcI6h/jI2Mtka6pt4Af3HKXhqDWa5NMEnn0yM9juHHaP8Lm1U8Cw2YFmMaQ6tlyU9o6WD3oHSQfdkKXePPZlI9pda+sB/WVbcm8hC3b5+911sNMmMCNruhdNHLeaPspfFyQr7923MYWDesQfoFSdv0mmM5olLneXIvpDgLFbAkpMJ7om5nG6IClFimoHziWFEz+q+cDbWNlr7Ved1OMtHZ2kPrwB5Fmcevz/ainn00D1O9LykKnExD3+CscbYTwd+Yg1OxvhPB46yQbShYA0bnaDzS4KdFHs1jjY59maWjlgKWZalF0dZDWEoym4jjMAuYUWUPUi4JsqmCKsxVsui/qViRKwttKbI/RytLRXUb9cN8ix1g1gvfMBisb7cwErE+nIDC4n15QYWFuvLDSwi1pcbcNv3yDjbQ+uEQVgs1os9L5WweJAwxOLFhGEWDxFGWDxCWMHiFT5xbrz0ItxjjWFjdc5QTq/vQV+v8BvcPNq3rHNGY+QDG51tMT1HnbwdfqnUXHjOt9uvc6C+26PTuqPsvPCcf7e/mAON3Z7iw5IX2Gvz8vd69B6pK2qxdcEWj94tdSUNtq64waN3SXvhBtteqMFT/HHJK2+2eZFmj94pdZWHbF3FIY/+KasPZ8iTEzHlp8MaG5jstpKTDH5De3Xax3DsJege07mXOAfhBzlJ36g7J6dXpb1C0WrFi3K6f7q8R6m5tues7QXMleHr8/XmPJljz1WIsRz4fifDeYTmpNSjizkq9es/s/onb6Q+sN5JitXfy+/l9/J7eWv5f/uDs5147jnl5FNB8Tw0GNMHck/yJB/N5e5eWg9ynHzMzbXJX2T3T+boLJbISeJ8Z0445DMq9kHo6DkV9jU1HOs5+/Sd2EJ4Llnxj2CzL/YtzIU1VuFjCuf3KFzpdSv8Hp6jjLmVvl7e6+pVcgYg2qSL835VUfptHV1r/TnVxye8aJc55LNazqfyCb2jvZrFx7BQij7w3qTYw3iqT038evuRmM/vG/P7fPf4ff5e9CPhcjqT+qQ+rLNro2frWNiqU8Qe6gxjIXqeF3aw6BlVDfWxSHgax1iOM3PYUd7ap1aEp+kszfpn85ZjN5VV2cqwnvhdDoVDVwDdWcbD9Hww36Vhl+eeZU6NzT0jxO8scdNLjMU5ocLizKqzz3a0b7TqsM440Q+PJ0Rrk+L1RM/4fKE+z5rwtMdTrub4ooq3qrXPVx2e9vnKFbRZiDbVnA2ttB8NuHIVFbply217tNYFXcsUaluhfsCmW+yl7H6gfZUQe1TFqjPEvvSs4Uc/sJ/C3vKM3x+iZ8LuPH/0TCAQ6vNvCk/7/eWewkDUnbeltS+wNTwdCJS70eZytOkpvKLVjTaLXKWKB7oVy217dOYodq1WqG039QM2ad31yH6gfQ8h1l+3VYe9mMrE2cOqw34sC/0oKgqJc8jyouiZkpJQX9GV4emionLfypKovvx9rX0lNeHpkpJyHW2uRJu+lVe30lmmxLVe8UFXuty2R+eakGubIvZj1A/YpLO+T/ZDF2f3m+jsrlt1IXG+on3dmXA4ZKwKR/vCdeHpcLjcWPWBVgN2Vy1n4llE2FWrEE88N5C2YMM4ir2LpWwvt+KMlbFnZ3jpVTzyB9tLImu3xyx1h9BfBr2T9MycjTgcrPwjO0oi67bHytfuiAnd9ZzRfsxi83wG/mI8i89zHEtwaB9n6fM8dSlbvnlO1hIc2v9Z7nmespQtzzzHtQSH9o1W/jzPs5StwDzHuwSH9puWMc9zL2XLP8/RluDQPtUqnuf5lrJVMs/JWYJD+1srOM/Tl7JVNM/JXoJD+2IrNM8zlrIVnufkLsGJXA5/3A4/RfI0/P57KG/+T/6hWOQxZpG/OlLXH864jjgWXldkXF+dcX2Nw2oNJSdyk8qoYim9WDQYlSsoo/NsGfJ5jF8aZA51ij4X9pzNO5VzNv9Ugt8YT/DtsfnzbHLCjTOeBR0TnyNeehFVI1Qu4ov93Uzk+gcUssWxvtHa4fcqY686khNajm/MwNrj9t3Vm5fPh/y03niHIFhz/K5+N9YgV69bXLtdSj/WrCGFJ3s9mjJWw/u+xLPzxrBOJbISrmSepY5R84qqJCIuPeaCTj2ljqqW2kvfEFOtgbFCJBE1OdCn3pvSI+laxgDmWIwpwupmeSmtbegufmf4ndfUzTqEjscc0Cnid3eshrtxrq+bdeNaAUcrpc9DsRblYnyX+Ssj0DmF7lczGnIwzak7YnMKwMm/3F9JZ+758rpZtZRmBzxOn5FjPQEvsM6oZLCVRX0B1y3WsqIX7E9ssb6AE9yoVDrAcYm+4dxFn4lhHXTCpkfatJ+PY20Av+g9vNIJvleU/WrGAa7gOG3OCnCK38sr6VlAqtwFW76ULcXmrQSv5P280gVbOaJtfY6vgq+n+C6bXwp+6BpeqYKfLfi+OX4W+EaKr9r8VeCHP8grs8DPTevrWeZcJXi0HnFHKIE9Qh+7bRXtDxKOfQd7b3NEhpg76oCNx2CDNboqneqmWdaY9ZTSqFUOM0/0FNOjn2S+6GlmRO9jfvtzRJcLc/arGSdQPFMCuhq9lWpjdmVWY06luzG3UmvMq0z5vgscFzi0tqvALGAW0C3uk0vRxL11KXSGxf0aIcQUHyeEcogQUzNMiCEPEGJIg4RY7voIsYL1P8iU6GeYGn0IY+I0lsasShrHwxjHBYzjMYzjyxjHV37DcfyG/R+T/f+U7P8p2f9Pyv6Pyv6flP0/Ift/1xT6b6H/tfQsEmOg7xmkj8PCOK4TZXp0p0Bf9AaBRnSXQH90N7P9/10a2/1ybH8sx3Zaju0+ObZ75NjulWO7W47t0xbGdpvooxptkOOLLjK+Fjm+Tjm+j8vxHZXju+PdHd9n5fjOyfE9KMf3GTm+B+T4zsrxnZHj+xMaX0KOr1eOb0ggiw7LcbJGR+Vc3+WYKVbZImOR/fkz2Z/Py/58jhDm/vQBEc88muTFYYuz8l9aVs5ZzsOI8dYznOMcUBK5Hxyfk0cp7qc5L0fghHHdOlUSjuZauWPkmSPsWNkAj5V9gX1zJsdRy/vDPORh1kzuMUt18kDs7E1/tMpC3UtZNO/Y/zoeVkd4JMJdasxTpkS9TjV2oWTnZVTGHZ9RrZJyMQ9O9ssJ4vSXKyH6zgjVt0oqo7SwjSgVkf5KNeRZ7UZ9T4w+6xP2LY9Ym2if6sul7zjhDACbI+qaiMUGJ6gHHjYwYaz8RAsz6l/+b8yRoHYU6sta21aqjlVSJc5HuaUJwZ1K4/ZXuUP0Wajdpw3RGvBG3NWR/g2ekGe9Dju+mIWzC5WL57jUB8+mSF5pUth6UtpSqd3NNj/Fs0q20HyzgtI+wX0qjdu/RQ9Zui7bvSI6Re3qWyP9V/hCnm0G7PhjFs5IVC7OHdSu78pIYemAsPU1actN7V5l81M8q8QvzmTLSocE9+k0br/fCFmGIeyOGO+LeN6fHxVnGdQ7VZwfXV56QtR5RtbRUedsfjB8JhgMeZblR/vy71k1nZ9f7ivIb11Rekpwn01x4a9/U3AtG8g/XcYU5wXyHf2+HKZLX/E6i2K68JWgaFNHm1rpmLDxzbT2PKuDUeJaIZxhIoyluMbKuwX3W+nctTZ3jrPqfsH5TjpnfQbnsjOC8910zuYMzuUPCM7fpnO2ZXDWPSg40+mcqxZyaH6NjZ8VvJl0npxz4z2fE2WzKDPe+2ctVE7Xz9F13kMtHnBTuu9Bxx+8uzc/r6A3r8Azad6WFznFTl7+IDt++eC+go9+ifXjnFVYTo/4pwI9Lb35gfLHCgM7enkw9mYguEtXVfebwWBUU1WdzpN0LqUzqVVUXP6GZTktnIUv0X/RWBFsRb0Pm7w4NlBUFHkzEIj4UNfCmZU+hx8IBiNP4F4/ll2mfAX3OLtwg2Idvb2FiQXOhf7svtzq+ngLnT+tro4W+lDR6jrUwoKEB1voWbTV1dLCKggPtLBqwuYW+o5lQuFT9FXLAeyPE/xSzOLfj9P3vuLOX8bizv8J+S/If0L+HfJzyBuQ1yAXIT+F/Cvkv0P+CfKPkJchP4S8BPl7yAuQ70O+B5mFTEO+C/k25FuQZyHPQJ6GPAWZgkxAnoD8NeTLkL+C/AXkzyEXIOch45AvQB6C/Bnkc5DPQh6EPAA5A7kfchpyD+RuyBjkFGQUcgIyAhmCDED6IElIAhKDHIN8AnIUcgTSA+mCdELaIW2QVkgUsh/SBNkHuQ3yMcgeyC2QXZCPQm6C3AjZCfkwZAfkQxCcU5wfgNRBrobUQN4HuRJyBWQrZAtkE2QDpBpSBVkDqYRUQMohEUgYEoKUQIohRZAgJADJh/ghBsQH0SEeiBuiQhQIhzC6jp9UkpMp33A4AlHHvV9sMehrLpx7aG/p53/ttZjT9n3e0/KNw+ad3BGM8nsfaUmcdLKv4zrJnyjrpc/RjjtZkk+U9ebz8secfMebXNnFKTYURcSG01EUdd77aAs996LnaPT8THEUR5V7H2uhz11FrNA6SGu5prTCBmLFHRtQ1cibnEcUihX5nZUBRYm86hiYcDlKoqojFM1yhKNurMHp8eN0ehScQeKuex9vUe/9y5ase7/U4r73K4gn9UIqrv5P4ykVPwl2z4TFjot1cyQwVTYSfKpspOjpspHiZ8pGSp4tGwl9q2wk/O2yEfbdMovdO0F1LHYa+MunLD5NNq6e/Ohs3GL3URnwfuDL11j8eyh7+WsWW34hwc5MmLhfJu6biftn4j6auJ8m7quJ+2viPpu43ybuu4n7b8IPTPiDCb8w4R8m/MSEv5jwGxP+Y8KPTPiTCb8y4V8m/MyEv5nwOxP+Z8IPTfijCb804Z8m/NSEv5rwWxP+a8KPTfizCb824d8m/NyEv5vwexP+byIOTMSDibgwER8m4sREvJiIGxPxYyKOTMSTibgyEV8m4sxEvJmIOxPxZyIOTcSjibg0EZ8m4tREvJqIWxPxayKOTcSzibg2Ed8m4txEvJuIexPxbyIPmMgHJvKCifxgIk+YyBcm8oaJ/GEij5jIJybyion8YiLPmMg3JvKOifxjIg+ZyEcm8pKJ/GQiT5nIVybylon8ZSKPmchnJvKaifxmIs+ZyHcm8p6J/GciD5rIhybyoon8aCJPmsiXJvKmifxpIo+ayKcm8qqJ/Goiz5rItybyron8ayIPm8jHJvKyifxsOi/Finv5QODBQC/vdQ4FLLfYwwUs7R7xnMHtTrDt5pdcmpbMdTp7x8n3sbfM4rfEz2M9HMceQeimLPW864qYrfuMrXvKUldgHfw1/dOWau//eGLR8mfs8qmlyp+1y59cqvxbdvlTS5V/2y7/2lLl37XLn160nPZL2Yx44/fRw1jad2XyHha8v8vgPbuEve9n8L65BO/5DN63luC9kMH7zhK8FzN4312C9/cZvL9dgveDDN70EryX0njjUs+2J79k77fsOvb+iieo7oq5/RVPjC/YW/EEfW70qlOJu48+39L1EXpecfGFVxkbD9l4kb5rMM5+8dx4gRNHHKF7LaW7mPu4GrJ1F4t7i3tT+teyH8f9z+Bq4BoLbb5WBl4wg1cKXjCDtwK8UAZvGXihDN5m8CoyeOvBq8jgrQWvOoO3GrzqDN614G3N4NWCtzWDdxV4NRm8beDVZPCuf1yldVuUX4c6il0esvEN4l2kMpprXeh+ntK9kft4VkhfyPt59uNZLD+Dp4GXn8ErA684g1cKXnEGbwV4kQzeMvAiGbzN4K3J4K0Hb00Gby14mzJ4q8HblMG7FrwrM3i14F2ZwbsKvLoM3jbwbN0b41993jvHvf7xLJprwbnu8axHvvqK91X23Myr66ayzuP3R7/6Y/v6lW9kcZaYvAjdz9J0r+H69bRr+k7vD19FX2DzjVdeKXz1lecLSfejBbpXCsfZ38yMf/Vrg9au5730XbDg94fF98MS0887oLuddK/W7oonZp93/Bsbmh1/3smoziOizq12nVtSdW5FnVtFnUdEnVsdP6c6tzrZI6hznurcsteus2f42M9/vNvJRL29DuhFvfO1t6DeXscvUO+RvXa9RzPqjf/L0vX+Pa3excx6P1m63sW0ej/LrPfi0vV+llbvtcx6P1i63mtp9V7PrPcPS9dLqOax12Xd8R8RD/d5nD0i4vRW+/pHaddJ162Fj+DeJ593uSzFvl+c1U0Hb3GK76s/Qt+bm/YylMn7psSFbtbL8gK8lxdok0m0l8Q9/Llyq+MNYWe+r45UP5VUPxXRT1EHfTxf89hz4y5nMPnj3d7z9VnO5D/v9Sb/pd6b/Mkr3uSLL3qTP/iBN/kP/+BN/sjlbXV8aHp8I2Ovs+/PjG9+KIvm4PXdu3Gg/mH8dfbZ2XEMjzjnwRkH57zkjO8lzh7sT+Y5j0rOoylOPXE+Fn80jfOG5LyR4rxCnJfjb0jO0J7dN4+7Ns6+vrvjNPEvgj/08t6bz0P36N6O06f27C17c8/e1eOqopKdi7ADX/m1dn4m2/kZyhP8hTj8Av42X/6aLH9NlP99HPcffjVfnpTlSZS7u15pGXe50Mau+Dg4SRf6+YLr5iT1ydVxemiX3eefoc+MPf3iG8jXp3btLntzV/3qJPpJ9+u8cf/Mow7GHsE6ysT1feI6hntFX5I8BjwGNOz6Yo04Rnl+7y9cIVs3Xi1tHwH3CLg9wB5gF7AL2AnsBFLdHlqP9z7nYm67Luk6hb2vuQxbd7Ed/Hbw24BtwFZgKzAKjEo7bcLO465Qmp2osHPOVS3t7Ad/P/hNwCbgPuA+4G3A26SdJmHnUy5mzNu5TdgxXYatu/gx8D8G/h7gHuAtwFuAu4C7pJ09wk63K5RmZ5ew0+iqlnY+Cv5Hwb8JeBPwRuCNwJ3AndLOTcLOH7pYcN7OTmHnGpdh6y5+GPwPg78DuAP4IeCHgNuB26WdHcLORlcozc52YWeVq1ra+QD4HwC/DlgHvBp4NbAGWCPt1Ak7eS4WmrdTI+w4XYatu/g+8N8H/pXAK4FXAK8AbgVulXauFHZ+oYTS7GwVdn6sVEs7W8DfAv4m4CbgBuAGYDWwWtrZZPuPwirm7VTb/qMYtu5iFfhV4K8BrgFWAiuBFcAKaWeN7T9KKM1Ohe0/SrW0Uw5+OfgRYAQYBoaBIWBI2onY/qOw6nk7Idt/FMPWXSwBvwT8YmAxsAhYBAwCg9JOse0/SijNTtD2H6Va2gmAHwBffGcfKL7DT7FJ3+mXdvJt/1HY1nk7hu0/imHrLor/CwC++O460AP0AN1At7Sj2/6jhNLsuG3/UaqlnU+A/wnwjwKPAlWgClSACpADOZABGfFoH7n3xy6ypdh+pdifMwp/Uvb8sm/yG8hDX3fuitN3sCoYY/lWYsxFHwZsT3wpl+X25h69tYW+i0B7faxND4VsPE/7e+ATZO8h2HsI+9CQrZug6/O558S+nri5R25vId0T2efEvp7qWSwhnqHEn7LrT2jnxP6e6odsnIp/zS47X3ZO7OnJVkr3ROk5sacnW7SPBz4Vf1raWnFO7OPJVko3teyc2MeT3ZCNT8e/Lu1vPif27sK+1D2x/pzYu0P3DO3Xgc/Gn7HLnlp7TuzXqU2al6/SPHxDtr36nJo8bN5psaR4NpSqM3XtOTWBtSOpuMrGUU7z9gS1Q/Nm2xih794R9+nac2pOL7f7dhX65hbl5+k+DkE3kvsLlfahIbdsG/oB4mavEnqad7JH+j5qQ1tu8239UEo/UpYn5jydO1SaJ85UxPPzUHyOuyJP3IMF3GV54h4ssLk5T5yrFvDW54n7sYC3Nk/cD+KFbDxBviD4q/PE/RB8qRu5Nk/cD+LT/QCOki9Q2YnaPHE/hA2pG7oqT9VDuV7DyPXSPD6z7Zw6Z2tbnkq6Z6+HTvZn9Hpb99R158S5S9i9Lk+1eFHc/v9TicmHcvO8D+cWeIexzx92VKhfwO9fzC20rwvWqX247k+7HsD1YNo17Q+Tw5gX2B8pKCgcLsgrTOwucCTzC8qobGSYHaey44WFhcOFeYXHsT99KPeyASuwTOz5RnH9MK55cPiYVbxS6Bj2k7zEyRJFyxzgiX3gSdpDFjnZKfC/AP6J2kB88qPBOOrcbv//qlWi7sna4ngiuNLBw8PHxsD9IrUVuUyUJYpWOcAT/E/WhsBb5eDlw8f6wOsD7xT2mKdWUbuXOVBH8D5VGxG8fnD6yZbHtvVJcMfAtdxl9h44MnxsAJwBcAaBg8BEUZkD5cLOWK0bdspEe5YvV9Tpg42+MtmeJ9WeR/COo+z4SicT4w46WT+u+8G19Pn2RqEbBWe0cKVTtKWn2tLttozhYydLVzpPrqLyXAfaFeVDtb74SQ1670rnGGwMUB+CuY6TZdnOk7gezHWygZU5Yu+fHGTJV5/PLhzIyy58NTvXZRn2+GkcdF8HXTnO4ew8lxiDYY9B3PdBNmT5C+x5D7ocPH/42Kdqy+F3tdMJlcZc4EC54I/U+tHfArHH/xT2n0mM+eeKkw3jergAWAj/XLXcK/Nr3GK98jl0WDyHHl4Bf9RWePuwpx1x1qm0xx1Buh8uXeYdLiv19hVo2N9G4mPY3/ZpTka845J33EGf0fMLw16vt6+QeIH48TTeCck7QfZWrvQmeH68bxnxgvETaby7JO+uRfa9fcuJXxS/S/KHAtrNfdhTj2lecQ4YRZ2hosKbj0N3otB7+lSgsOzNgLa6D/vr47A9CtvHVxTOtXm80MlI3y/b7HfQl0q+P3MS19TOca93bhz9XvKd78+cktxT4Gauuf2lXgeVnUTZXbLOUMB78+jKFY4RXI+ugA3qm9d7OrdrRUtqjz6UX3DzCPR3FaDPiPc3872r+9FnOieMiu/xJCb7cDag/jnE7/eJ3w27/ETqHDCi2eeAUZlfaY0/of3YVS15qbNAbueKltR5IHUOGNFwDpB1U2eEzL1/6nxwQsP5IK3tNlH/8bm2M88ImXv/1PnghHZurm+57StaUmeAEe1Tc33JPCdk7v9TZ4QTmrmgT3uEne65PmWeEzL3/6kzwgmtcb5PbXafbhK2/nCuT5lnhcwzQOqccEK7ZkGfdgg7G+f6lHlWyDwDpM4JJ7RV831qtftUJ2zlzfUp87yQeQ5InRVOaM4FfbrS9hsl1afM80LmOSB1VoBfKXN9itp92mT7kZLqU+aZIfMskDovwJ+U9D6tsf1prk+ZZ4bMs0DqvAB/mu/TfrtPEduf5vqUeW7IPA+kzgzwpwV9Krb9aa5PmeeGzPNA6swAf5rvU5Pdp3zbn+b6lHl2yDwTpM4N8KcFfdJtf5rrk2xnNPMM8XZnAdr7uG0/U1LnAfjW3HlgVHMqH3f0Tfru4qODuZozmeX18mRitC/Pe7vClITL4Upmo6z/Wnec9MnbXGUYXCJJn9PmaarL4UieVbJWUR7X/4sblNf1X3KD8qb+Bjcoj+r/zg3KT/q/coPyon6RG5SH9Je5QblK/yduUB7SX+AG5SX9JW5QftGnuUH5Rv8eNyi/6M9yg/KN/m1uUM7Qp7hBOUR/mhuUM/Qvc4NyiP4ENygX6Be4QblB/wtuUC7QH+IG5QZ9nBsU3/qD3KB41z/HDYpv/TQ3KN71M9yguNVPcYPiWL+bGxS3+hA3KI71E9ygWNQT3KDY1Pu4QbGoH+UGxaZ+jBsUY3onNyjm9B5uUIzpUW5QzOlt3KC40W/jBsWR3sQNiht9FzcojvQ93KB40Hdyg+JDv4kbFA/6dm5QfOg7uEE+rtdwg3xer+MG+bi+lRvk8/qV3CDf1au5Qb6sb+IG+a5ewQ3yZX0NN8gf9RA3yD/1CDeEPwa5Qf6pF3OD/Es3uEH+pudzg/xLd3OD/E3XuUH+pTNukL/pCjeSzD2bdDmYWM81p1us75pcP4FiPdXkMyugWB81+cxKk8+wNLkeaXJ90uQ6o8lnU5pcZzS57mhy7dDkMydNrh2aXEs0uSZo8lmSJtcETa4RmszzmnxGpMk8r8m8r8n8rclnP5rM35rM55rMyZp8pqPJnKzJHK3JXKvJZzWazLWazL2azJ+afAajyfypyXyqybyoyWcrmsyLmsyTmsx1mnxmoslcp8ncp8kcpslnIZrMYZrMaZrMS5p8xqHJvKTJPKXJPKPJZxeazDOazDuazDOazDvAU2z4TpY49lLyWJzVEB6NswRhT5xNEXZiO9oLbEM5YRTlhE0oJ7wN5UngHpQT7kI54U0oJ9yJchO4A+WE21FOWIdywhqU9wGvRDnhVpQTbkI5YTXK+4FrUE5YgXLCCMoJQygfABajnDCIcsJ8lBMaKB8E6igndKOcUEE5IYuz1lBiopgNTfr62GgC+/DsPj7Kkn2j9P8LHDjrZ7lcScql9rOCBLM4iwW7HnuSe4aPUW6lnJqwvxg/k6h5mHa+s/30Ir9r6Pv47lkQmNOlVAW71r9/wKXgiN/5EuPoF+udTL0zN6jZmHrdcDDjvXmGx0YJrFC+7+0/LlkdhA9pC9+P+4K2xH84rL2+LtRIb5EMyVcOhjZsqKquqj7Q0hOqiPb0dHZfsW4dLqKH91U1dhxad2NDU3NHu3gBZcehzsM9zeva2o4cWtvZ1dHa3NgT6upoXPueKlgIbdhSvX5T6D37tlZXN6zftmnLlg0btu3buuk92xres7Wpef229Zs2NjVu3Ni4sWHftn2bV7MdLe0Hm7uuCFGHduyok90IVaxr62g/0NxV39nQE6WXV9Z3d9T3RBt66rs6D3XXNza013cf7uzs6OoBo/Fgw4GW9gMob65vat53+EB9S/v+jvr9HV31DW1t9R3d9ejn/pa25u513V2NC3ouLn6D7rK3fy8jY/9mbc3QR6Q+k/9BqV+Xoe9Ygv/HS+ifkXpHhv4nS/ALHIvrr1lCf3AJ/b1S/x+ZBXxxfovUl2Ton1+Cf51zcf1fOxcf7yplcf6nl9BrrsX1ty+h/8kS+l8toc9VF9e/l15czd3z79WWP2Xq4u/1fL+6+Hs996iLv9fzCXXx93quzFr8vZ6sqr2jp5lVNd3R3n3HIVZ1oP1wVbShO8rkv6Tv6WJVXR1NDT0NrKqn+WiP0NJbUqFubuvqQCA2NSEMWdW+7m5WZb+Ou6qx264hL+nNqYcajrYcOnwIHGQWek0so1e/9jTsA0bRjP2bjW/xus/2w4fq6U2sb0NpeFvKvuYueknt25uiPNRUv6+hq6ul+a24aW/Btd99uzRVvG32SGPj21HEe327G7saehqjb8HF3aJ314pWm+y3174Nu6u58bBYAN6G19Le1AJuD1JvW9u7/h7UuXv7u2io4XfUUJqf/S6GtdBb3+0Wf83n3+0G5yLnd9LQgvh7t1tcLIp/F23O54LfRWsLMwq9IbvzsHiH+VzwZ+gaFtG9ddb+//pl0fV1uz9ce/111/zfeD+JQ/6tjdRryZf6+xwL3uGd9pMlzzQ845yTwlNp5xxH2nulU+efHMh/WlZHqn7qnJPC6vRt5yLvpV4mbfOMc1EKt2bUVzJwpXxfOc84h6Ww0LHE32GRP5dJXar+3D7Ou/D93Zl/7yb1sy79b8ek/50ZPePvzCzRgc2y7tw711P7zt0L/56LK2P+Uq8Rf5/Upep3pv4OjKx/Sn3r8V+fWX94YX3HW3ef3biILr0+e5v6Db9l/SNSV71E/W1L3P8UxtPuffrPhVtsnObz871sEf8fXOJ9+Ym9Nt73Nv3/5BL1l33Mxrqst67/vwEEsPdu"
   
- BLOCK_H = 16
- BLOCK_N = 64
- BLOCK_C = 128
- BLOCK_R = 64
- SPLITKV_BATCH_1 = 4
- SPLITKV_BATCH_2 = 32
- SPLITKV_BATCH_3 = 64
- SPLITKV_KV = 4096
+ _SRC_PATH = HERE / "_submission_mla.cu"
+ _source_text = zlib.decompress(base64.b64decode(MLA_CPP_B64)).decode("utf-8")
+ if not _SRC_PATH.exists() or _SRC_PATH.read_text() != _source_text:
+     _SRC_PATH.write_text(_source_text)
   
+ _BINDING_SRC = r"""
+ #include <torch/extension.h>
+ #include <vector>
   
- @triton.jit
- def _fp4_table(x):
-     x = x.to(tl.int32)
-     v = tl.where(x == 0, 0.0, 0.5)
-     v = tl.where(x == 2, 1.0, v)
-     v = tl.where(x == 3, 1.5, v)
-     v = tl.where(x == 4, 2.0, v)
-     v = tl.where(x == 5, 3.0, v)
-     v = tl.where(x == 6, 4.0, v)
-     v = tl.where(x == 7, 6.0, v)
-     return v.to(tl.float32)
+ std::vector<torch::Tensor> mla_decode_h16_fp8_fp8(torch::Tensor q,
+                                                   torch::Tensor kv,
+                                                   torch::Tensor kv_scale,
+                                                   int64_t splitkv,
+                                                   float softmax_scale,
+                                                   std::optional<torch::Tensor> hsaco_tensor);
   
+ PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+     m.def("mla_decode_h16_fp8_fp8", &mla_decode_h16_fp8_fp8);
+ }
+ """
   
- @triton.jit
- def _decode_e8m0(scale_u8):
-     return tl.exp2(scale_u8.to(tl.float32) - 127.0)
-  
-  
- @triton.jit
- def _load_mxfp4(
-     packed_ptr,
-     scale_ptr,
-     token_idx,
-     dim_offsets,
-     mask_n,
-     mask_d,
-     PACKED_DIM_CONST: tl.constexpr,
-     BLOCKS_CONST: tl.constexpr,
-     BLOCK_SIZE_CONST: tl.constexpr,
- ):
-     packed_idx = dim_offsets // 2
-     packed_ptrs = packed_ptr + token_idx[None, :] * PACKED_DIM_CONST + packed_idx[:, None]
-     packed = tl.load(
-         packed_ptrs,
-         mask=mask_d[:, None] & mask_n[None, :],
-         other=0,
-     ).to(tl.int32)
-  
-     lo = packed & 0xF
-     hi = (packed >> 4) & 0xF
-     nibbles = tl.where((dim_offsets[:, None] & 1) == 0, lo, hi)
-  
-     sign = tl.where((nibbles & 0x8) == 0, 1.0, -1.0)
-     mag = _fp4_table(nibbles & 0x7)
-  
-     block_idx = dim_offsets // BLOCK_SIZE_CONST
-     scale_ptrs = scale_ptr + token_idx[None, :] * BLOCKS_CONST + block_idx[:, None]
-     block_scale = _decode_e8m0(
-         tl.load(
-             scale_ptrs,
-             mask=mask_d[:, None] & mask_n[None, :],
-             other=0,
-         ).to(tl.int32)
+ def _build_inline_module():
+     return load_inline(
+         name="mla_standalone_ext",
+         cpp_sources=[_BINDING_SRC],
+         cuda_sources=[_source_text],
+         verbose=True,
+         extra_cflags=["-O3", "-std=c++20"],
+         extra_cuda_cflags=[
+             "-O3",
+             "-std=c++20",
+             "-U__HIP_NO_HALF_OPERATORS__",
+             "-U__HIP_NO_HALF_CONVERSIONS__",
+             "--save-temps",
+             "-ffast-math",
+             "-fno-finite-math-only",
+             # "-mllvm", "-amdgpu-mfma-vgpr-form",
+         ],
+         extra_ldflags=["-lamdhip64"],
      )
-     return sign * mag * block_scale
   
   
- @triton.jit
- def _decode_mxfp4_tile_with_scale(
-     packed_tile,
-     scale_t,
-     BLOCK_SIZE_CONST: tl.constexpr,
-     CHUNK_DIM_CONST: tl.constexpr,
- ):
-     offs_d = tl.arange(0, CHUNK_DIM_CONST)
-     packed_idx = offs_d // 2
-     packed = packed_tile[packed_idx, :].to(tl.int32)
-     lo = packed & 0xF
-     hi = (packed >> 4) & 0xF
-     nibbles = tl.where((offs_d[:, None] & 1) == 0, lo, hi)
+ def _load_hsaco_tensor():
+     if not GPU_HSACO_B64:
+         return torch.empty((0,), dtype=torch.uint8)
+     hsaco = zlib.decompress(base64.b64decode(GPU_HSACO_B64))
+     return torch.frombuffer(memoryview(bytearray(hsaco)), dtype=torch.uint8)
   
-     sign = tl.where((nibbles & 0x8) == 0, 1.0, -1.0)
-     mag = _fp4_table(nibbles & 0x7)
   
-     block_idx = offs_d // BLOCK_SIZE_CONST
-     block_scale = scale_t[block_idx, :]
-     return sign * mag * block_scale
+ _hsaco_tensor = _load_hsaco_tensor()
+ _module = _build_inline_module()
   
   
- @triton.jit
- def _mla_decode_kernel(
-     q_ptr,
-     kv_ptr,
-     kv_indptr_ptr,
-     o_ptr,
-     kv_scale_ptr,
-     NUM_HEADS_CONST: tl.constexpr,
-     KV_LORA_RANK_CONST: tl.constexpr,
-     QK_ROPE_HEAD_DIM_CONST: tl.constexpr,
-     QK_DIM_CONST: tl.constexpr,
-     V_DIM_CONST: tl.constexpr,
-     SM_SCALE_CONST: tl.constexpr,
-     MAX_KV_CONST: tl.constexpr,
-     BLOCK_H_CONST: tl.constexpr,
-     BLOCK_N_CONST: tl.constexpr,
-     BLOCK_C_CONST: tl.constexpr,
-     BLOCK_R_CONST: tl.constexpr,
- ):
-     pid_b = tl.program_id(0)
-     pid_h = tl.program_id(1)
+ def _quantize_fp8(x: torch.Tensor) -> torch.Tensor:
+     return x.to(torch.float8_e4m3fn).contiguous()
   
-     offs_h = pid_h * BLOCK_H_CONST + tl.arange(0, BLOCK_H_CONST)
-     offs_c = tl.arange(0, BLOCK_C_CONST)
-     offs_r = tl.arange(0, BLOCK_R_CONST)
-     offs_v = tl.arange(0, V_DIM_CONST)
   
-     mask_h = offs_h < NUM_HEADS_CONST
-     mask_c = offs_c < KV_LORA_RANK_CONST
-     mask_r = offs_r < QK_ROPE_HEAD_DIM_CONST
-     mask_v = offs_v < V_DIM_CONST
+ def _dyn_quantize_fp8(tensor: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
+     finfo = torch.finfo(torch.float8_e4m3fn)
+     amax = tensor.abs().amax().clamp(min=1e-12)
+     scale = amax / finfo.max
+     fp8_tensor = (tensor / scale).clamp(min=finfo.min, max=finfo.max).to(torch.float8_e4m3fn)
+     return fp8_tensor, scale.to(torch.float32).reshape(1)
   
-     q_nope_ptrs = (
-         q_ptr
-         + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * QK_DIM_CONST + offs_c[None, :])
-     )
-     q_pe_ptrs = (
-         q_ptr
-         + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * QK_DIM_CONST + KV_LORA_RANK_CONST + offs_r[None, :])
-     )
-     q_nope_bf16 = tl.load(q_nope_ptrs, mask=mask_h[:, None] & mask_c[None, :], other=0.0)
-     q_pe_bf16 = tl.load(q_pe_ptrs, mask=mask_h[:, None] & mask_r[None, :], other=0.0)
-     q_amax = tl.maximum(tl.max(tl.abs(q_nope_bf16), axis=1), tl.max(tl.abs(q_pe_bf16), axis=1))
-     q_scale = tl.maximum(q_amax / 448.0, 1e-12)
-     q_nope = (q_nope_bf16 / q_scale[:, None]).to(tl.float8e4nv)
-     q_pe = (q_pe_bf16 / q_scale[:, None]).to(tl.float8e4nv)
   
-     kv_start = tl.load(kv_indptr_ptr + pid_b)
-     kv_end = tl.load(kv_indptr_ptr + pid_b + 1)
-     logits_scale = q_scale[:, None] * tl.load(kv_scale_ptr) * SM_SCALE_CONST
-     v_scale = tl.load(kv_scale_ptr)
+ def _choose_splitkv(bs: int, seq: int) -> int:
+     NUM_CUS = 256
+     NUM_WARPS = 4
+     TILE_K = 32
+     # (bs * split) % (cu * warp) == 0
+     # splitk <= 32
+     # seq % splitk == 0 && seq / splitk >= 32
+     MAPS = {
+         (4, 1024): 32,
+         (4, 8192): 32,
+         (32, 1024): 32,
+         (32, 8192): 32,
+         (64, 1024): 16,
+         (64, 8192): 16,
+         (256, 1024): 4,
+         (256, 8192): 4,
+     }
+     override = os.getenv("MLA_SPLITKV_OVERRIDE")
+     if override is not None:
+         return int(override)
+     mapped = MAPS.get((bs, seq))
+     if mapped is not None:
+         return mapped
+     split = min(NUM_WARPS * NUM_CUS // bs, seq // TILE_K, 32)
+     return split
   
-     e_max = tl.full((BLOCK_H_CONST,), float("-inf"), tl.float32)
-     e_sum = tl.zeros((BLOCK_H_CONST,), tl.float32)
-     acc = tl.zeros((BLOCK_H_CONST, V_DIM_CONST), tl.float32)
   
-     for start_n in range(0, MAX_KV_CONST, BLOCK_N_CONST):
-         offs_n = start_n + tl.arange(0, BLOCK_N_CONST)
-         token_idx = kv_start + offs_n
-         mask_n = token_idx < kv_end
-  
-         kv_nope_ptrs = kv_ptr + token_idx[None, :] * QK_DIM_CONST + offs_c[:, None]
-         kv_pe_ptrs = (
-             kv_ptr
-             + token_idx[None, :] * QK_DIM_CONST
-             + KV_LORA_RANK_CONST
-             + offs_r[:, None]
-         )
-         v_ptrs = kv_ptr + token_idx[:, None] * QK_DIM_CONST + offs_v[None, :]
-  
-         kv_nope = tl.load(
-             kv_nope_ptrs,
-             mask=mask_c[:, None] & mask_n[None, :],
-             other=0.0,
-         )
-         kv_pe = tl.load(
-             kv_pe_ptrs,
-             mask=mask_r[:, None] & mask_n[None, :],
-             other=0.0,
-         )
-         v = tl.load(
-             v_ptrs,
-             mask=mask_n[:, None] & mask_v[None, :],
-             other=0.0,
-         )
-  
-         qk = tl.dot(q_nope, kv_nope) + tl.dot(q_pe, kv_pe)
-         qk = qk * logits_scale
-         qk = tl.where(mask_h[:, None] & mask_n[None, :], qk, float("-inf"))
-  
-         n_e_max = tl.maximum(tl.max(qk, axis=1), e_max)
-         re_scale = tl.math.exp2((e_max - n_e_max) * 1.4426950408889634)
-         p = tl.math.exp2((qk - n_e_max[:, None]) * 1.4426950408889634)
-  
-         acc *= re_scale[:, None]
-         acc += tl.dot(p.to(v.dtype), v)
-         e_sum = e_sum * re_scale + tl.sum(p, axis=1)
-         e_max = n_e_max
-  
-     out_ptrs = (
-         o_ptr
-         + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * V_DIM_CONST + offs_v[None, :])
-     )
-     tl.store(
-         out_ptrs,
-         (acc * v_scale) / e_sum[:, None],
-         mask=mask_h[:, None] & mask_v[None, :],
-     )
-  
-  
- @triton.jit
- def _mla_decode_splitkv_kernel(
-     q_ptr,
-     kv_ptr,
-     kv_indptr_ptr,
-     partial_acc_ptr,
-     partial_max_ptr,
-     partial_sum_ptr,
-     kv_scale_ptr,
-     NUM_SPLITS_CONST: tl.constexpr,
-     NUM_HEADS_CONST: tl.constexpr,
-     KV_LORA_RANK_CONST: tl.constexpr,
-     QK_ROPE_HEAD_DIM_CONST: tl.constexpr,
-     QK_DIM_CONST: tl.constexpr,
-     V_DIM_CONST: tl.constexpr,
-     SM_SCALE_CONST: tl.constexpr,
-     MAX_SPLIT_KV_CONST: tl.constexpr,
-     BLOCK_H_CONST: tl.constexpr,
-     BLOCK_N_CONST: tl.constexpr,
-     BLOCK_C_CONST: tl.constexpr,
-     BLOCK_R_CONST: tl.constexpr,
- ):
-     pid_b = tl.program_id(0)
-     pid_h = tl.program_id(1)
-     pid_s = tl.program_id(2)
-  
-     offs_h = pid_h * BLOCK_H_CONST + tl.arange(0, BLOCK_H_CONST)
-     offs_c = tl.arange(0, BLOCK_C_CONST)
-     offs_r = tl.arange(0, BLOCK_R_CONST)
-     offs_v = tl.arange(0, V_DIM_CONST)
-  
-     mask_h = offs_h < NUM_HEADS_CONST
-     mask_c = offs_c < KV_LORA_RANK_CONST
-     mask_r = offs_r < QK_ROPE_HEAD_DIM_CONST
-     mask_v = offs_v < V_DIM_CONST
-  
-     q_nope_ptrs = (
-         q_ptr
-         + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * QK_DIM_CONST + offs_c[None, :])
-     )
-     q_pe_ptrs = (
-         q_ptr
-         + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * QK_DIM_CONST + KV_LORA_RANK_CONST + offs_r[None, :])
-     )
-     q_nope_bf16 = tl.load(q_nope_ptrs, mask=mask_h[:, None] & mask_c[None, :], other=0.0)
-     q_pe_bf16 = tl.load(q_pe_ptrs, mask=mask_h[:, None] & mask_r[None, :], other=0.0)
-     q_amax = tl.maximum(tl.max(tl.abs(q_nope_bf16), axis=1), tl.max(tl.abs(q_pe_bf16), axis=1))
-     q_scale = tl.maximum(q_amax / 448.0, 1e-12)
-     q_nope = (q_nope_bf16 / q_scale[:, None]).to(tl.float8e4nv)
-     q_pe = (q_pe_bf16 / q_scale[:, None]).to(tl.float8e4nv)
-  
-     kv_start = tl.load(kv_indptr_ptr + pid_b)
-     kv_end = tl.load(kv_indptr_ptr + pid_b + 1)
-     kv_len = kv_end - kv_start
-     split_start = kv_start + (kv_len * pid_s) // NUM_SPLITS_CONST
-     split_end = kv_start + (kv_len * (pid_s + 1)) // NUM_SPLITS_CONST
-  
-     logits_scale = q_scale[:, None] * tl.load(kv_scale_ptr) * SM_SCALE_CONST
-     e_max = tl.full((BLOCK_H_CONST,), float("-inf"), tl.float32)
-     e_sum = tl.zeros((BLOCK_H_CONST,), tl.float32)
-     acc = tl.zeros((BLOCK_H_CONST, V_DIM_CONST), tl.float32)
-  
-     for start_n in range(0, MAX_SPLIT_KV_CONST, BLOCK_N_CONST):
-         offs_n = start_n + tl.arange(0, BLOCK_N_CONST)
-         token_idx = split_start + offs_n
-         mask_n = token_idx < split_end
-  
-         kv_nope_ptrs = kv_ptr + token_idx[None, :] * QK_DIM_CONST + offs_c[:, None]
-         kv_pe_ptrs = (
-             kv_ptr
-             + token_idx[None, :] * QK_DIM_CONST
-             + KV_LORA_RANK_CONST
-             + offs_r[:, None]
-         )
-         v_ptrs = kv_ptr + token_idx[:, None] * QK_DIM_CONST + offs_v[None, :]
-  
-         kv_nope = tl.load(
-             kv_nope_ptrs,
-             mask=mask_c[:, None] & mask_n[None, :],
-             other=0.0,
-         )
-         kv_pe = tl.load(
-             kv_pe_ptrs,
-             mask=mask_r[:, None] & mask_n[None, :],
-             other=0.0,
-         )
-         v = tl.load(
-             v_ptrs,
-             mask=mask_n[:, None] & mask_v[None, :],
-             other=0.0,
-         )
-  
-         qk = tl.dot(q_nope, kv_nope) + tl.dot(q_pe, kv_pe)
-         qk = qk * logits_scale
-         qk = tl.where(mask_h[:, None] & mask_n[None, :], qk, float("-inf"))
-  
-         n_e_max = tl.maximum(tl.max(qk, axis=1), e_max)
-         re_scale = tl.math.exp2((e_max - n_e_max) * 1.4426950408889634)
-         p = tl.math.exp2((qk - n_e_max[:, None]) * 1.4426950408889634)
-  
-         acc *= re_scale[:, None]
-         acc += tl.dot(p.to(v.dtype), v)
-         e_sum = e_sum * re_scale + tl.sum(p, axis=1)
-         e_max = n_e_max
-  
-     acc_ptrs = (
-         partial_acc_ptr
-         + ((((pid_b * NUM_SPLITS_CONST + pid_s) * NUM_HEADS_CONST) + offs_h[:, None]) * V_DIM_CONST + offs_v[None, :])
-     )
-     max_ptrs = partial_max_ptr + (((pid_b * NUM_SPLITS_CONST + pid_s) * NUM_HEADS_CONST) + offs_h)
-     sum_ptrs = partial_sum_ptr + (((pid_b * NUM_SPLITS_CONST + pid_s) * NUM_HEADS_CONST) + offs_h)
-     tl.store(acc_ptrs, acc, mask=mask_h[:, None] & mask_v[None, :])
-     tl.store(max_ptrs, e_max, mask=mask_h)
-     tl.store(sum_ptrs, e_sum, mask=mask_h)
-  
-  
- @triton.jit
- def _mla_reduce_splitkv_kernel(
-     partial_acc_ptr,
-     partial_max_ptr,
-     partial_sum_ptr,
-     o_ptr,
-     kv_scale_ptr,
-     NUM_SPLITS_CONST: tl.constexpr,
-     NUM_HEADS_CONST: tl.constexpr,
-     V_DIM_CONST: tl.constexpr,
-     BLOCK_H_CONST: tl.constexpr,
- ):
-     pid_b = tl.program_id(0)
-     pid_h = tl.program_id(1)
-  
-     offs_h = pid_h * BLOCK_H_CONST + tl.arange(0, BLOCK_H_CONST)
-     offs_v = tl.arange(0, V_DIM_CONST)
-     mask_h = offs_h < NUM_HEADS_CONST
-     mask_v = offs_v < V_DIM_CONST
-  
-     global_max = tl.full((BLOCK_H_CONST,), float("-inf"), tl.float32)
-     global_sum = tl.zeros((BLOCK_H_CONST,), tl.float32)
-  
-     for split_idx in range(0, NUM_SPLITS_CONST):
-         max_ptrs = partial_max_ptr + (((pid_b * NUM_SPLITS_CONST + split_idx) * NUM_HEADS_CONST) + offs_h)
-         sum_ptrs = partial_sum_ptr + (((pid_b * NUM_SPLITS_CONST + split_idx) * NUM_HEADS_CONST) + offs_h)
-         split_max = tl.load(max_ptrs, mask=mask_h, other=float("-inf"))
-         split_sum = tl.load(sum_ptrs, mask=mask_h, other=0.0)
-         next_max = tl.maximum(global_max, split_max)
-         global_sum = global_sum * tl.math.exp2((global_max - next_max) * 1.4426950408889634)
-         global_sum += split_sum * tl.math.exp2((split_max - next_max) * 1.4426950408889634)
-         global_max = next_max
-  
-     acc = tl.zeros((BLOCK_H_CONST, V_DIM_CONST), tl.float32)
-     for split_idx in range(0, NUM_SPLITS_CONST):
-         max_ptrs = partial_max_ptr + (((pid_b * NUM_SPLITS_CONST + split_idx) * NUM_HEADS_CONST) + offs_h)
-         acc_ptrs = (
-             partial_acc_ptr
-             + ((((pid_b * NUM_SPLITS_CONST + split_idx) * NUM_HEADS_CONST) + offs_h[:, None]) * V_DIM_CONST + offs_v[None, :])
-         )
-         split_max = tl.load(max_ptrs, mask=mask_h, other=float("-inf"))
-         scale = tl.math.exp2((split_max - global_max) * 1.4426950408889634)
-         split_acc = tl.load(acc_ptrs, mask=mask_h[:, None] & mask_v[None, :], other=0.0)
-         acc += split_acc * scale[:, None]
-  
-     v_scale = tl.load(kv_scale_ptr)
-     out_ptrs = (
-         o_ptr
-         + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * V_DIM_CONST + offs_v[None, :])
-     )
-     tl.store(
-         out_ptrs,
-         (acc * v_scale) / global_sum[:, None],
-         mask=mask_h[:, None] & mask_v[None, :],
-     )
-  
-  
- @triton.jit
- def _mla_decode_mxfp4_kernel(
-     q_ptr,
-     kv_packed_ptr,
-     kv_scale_ptr,
-     kv_fp8_scale_ptr,
-     kv_indptr_ptr,
-     o_ptr,
-     NUM_HEADS_CONST: tl.constexpr,
-     KV_LORA_RANK_CONST: tl.constexpr,
-     QK_ROPE_HEAD_DIM_CONST: tl.constexpr,
-     QK_DIM_CONST: tl.constexpr,
-     V_DIM_CONST: tl.constexpr,
-     SM_SCALE_CONST: tl.constexpr,
-     MAX_KV_CONST: tl.constexpr,
-     BLOCK_H_CONST: tl.constexpr,
-     BLOCK_N_CONST: tl.constexpr,
-     BLOCK_C_CONST: tl.constexpr,
-     BLOCK_R_CONST: tl.constexpr,
-     PACKED_DIM_CONST: tl.constexpr,
-     BLOCKS_CONST: tl.constexpr,
-     BLOCK_SIZE_CONST: tl.constexpr,
- ):
-     pid_b = tl.program_id(0)
-     pid_h = tl.program_id(1)
-  
-     offs_h = pid_h * BLOCK_H_CONST + tl.arange(0, BLOCK_H_CONST)
-     offs_v = tl.arange(0, V_DIM_CONST)
-  
-     PACKED_V_CONST: tl.constexpr = KV_LORA_RANK_CONST // 2
-     SCALE_V_CONST: tl.constexpr = KV_LORA_RANK_CONST // BLOCK_SIZE_CONST
-     PACKED_R_CONST: tl.constexpr = BLOCK_R_CONST // 2
-     SCALE_R_CONST: tl.constexpr = BLOCK_R_CONST // BLOCK_SIZE_CONST
-  
-     mask_h = offs_h < NUM_HEADS_CONST
-     mask_v = offs_v < V_DIM_CONST
-  
-     offs_r = tl.arange(0, BLOCK_R_CONST) + KV_LORA_RANK_CONST
-     mask_r = offs_r < QK_DIM_CONST
-     offs_c = tl.arange(0, KV_LORA_RANK_CONST)
-     q_nope_ptrs = q_ptr + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * QK_DIM_CONST + offs_c[None, :])
-     q_nope_bf16 = tl.load(q_nope_ptrs, mask=mask_h[:, None] & (offs_c[None, :] < KV_LORA_RANK_CONST), other=0.0)
-     q_nope_descale = tl.full((BLOCK_H_CONST, SCALE_V_CONST), 127, dtype=tl.uint8)
-     q_pe_ptrs = q_ptr + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * QK_DIM_CONST + offs_r[None, :])
-     q_pe_bf16 = tl.load(q_pe_ptrs, mask=mask_h[:, None] & mask_r[None, :], other=0.0)
-     q_amax = tl.maximum(tl.max(tl.abs(q_nope_bf16), axis=1), tl.max(tl.abs(q_pe_bf16), axis=1))
-     q_scale = tl.maximum(q_amax / 448.0, 1e-12)
-     q_nope = (q_nope_bf16 / q_scale[:, None]).to(tl.float8e4nv)
-     q_pe = (q_pe_bf16 / q_scale[:, None]).to(tl.float8e4nv)
-     q_pe_descale = tl.full((BLOCK_H_CONST, BLOCK_R_CONST // BLOCK_SIZE_CONST), 127, dtype=tl.uint8)
-  
-     kv_start = tl.load(kv_indptr_ptr + pid_b)
-     kv_end = tl.load(kv_indptr_ptr + pid_b + 1)
-     logits_scale = q_scale[:, None] * (SM_SCALE_CONST * 1.4426950408889634)
-     v_scale = tl.load(kv_fp8_scale_ptr)
-  
-     e_max = tl.full((BLOCK_H_CONST,), float("-inf"), tl.float32)
-     e_sum = tl.zeros((BLOCK_H_CONST,), tl.float32)
-     acc = tl.zeros((BLOCK_H_CONST, V_DIM_CONST), tl.float32)
-  
-     for start_n in range(0, MAX_KV_CONST, BLOCK_N_CONST):
-         offs_n = start_n + tl.arange(0, BLOCK_N_CONST)
-         token_idx = kv_start + offs_n
-         mask_n = token_idx < kv_end
-  
-         offs_lp = tl.arange(0, PACKED_V_CONST)
-         latent_ptrs = kv_packed_ptr + token_idx[None, :] * PACKED_DIM_CONST + offs_lp[:, None]
-         latent_packed = tl.load(
-             latent_ptrs,
-             mask=(offs_lp[:, None] < PACKED_V_CONST) & mask_n[None, :],
-             other=0,
-         )
-  
-         offs_ls = tl.arange(0, SCALE_V_CONST)
-         latent_scale_ptrs = kv_scale_ptr + token_idx[:, None] * BLOCKS_CONST + offs_ls[None, :]
-         latent_scale = tl.load(
-             latent_scale_ptrs,
-             mask=mask_n[:, None] & (offs_ls[None, :] < SCALE_V_CONST),
-             other=0,
-         )
-  
-         offs_rp = tl.arange(0, PACKED_R_CONST)
-         rope_ptrs = kv_packed_ptr + token_idx[None, :] * PACKED_DIM_CONST + ((KV_LORA_RANK_CONST // 2) + offs_rp)[:, None]
-         kv_pe = tl.load(
-             rope_ptrs,
-             mask=(offs_rp[:, None] < PACKED_R_CONST) & mask_n[None, :],
-             other=0,
-         )
-         offs_rs = (KV_LORA_RANK_CONST // BLOCK_SIZE_CONST) + tl.arange(0, SCALE_R_CONST)
-         rope_scale_ptrs = kv_scale_ptr + token_idx[:, None] * BLOCKS_CONST + offs_rs[None, :]
-         kv_pe_scale = tl.load(
-             rope_scale_ptrs,
-             mask=mask_n[:, None] & (offs_rs[None, :] < BLOCKS_CONST),
-             other=0,
-         )
-  
-         qk = tl.zeros((BLOCK_H_CONST, BLOCK_N_CONST), dtype=tl.float32)
-         qk = tl.dot_scaled(
-             q_nope,
-             q_nope_descale,
-             "e4m3",
-             latent_packed,
-             latent_scale,
-             "e2m1",
-             fast_math=True,
-             acc=qk,
-         )
-         qk = tl.dot_scaled(
-             q_pe,
-             q_pe_descale,
-             "e4m3",
-             kv_pe,
-             kv_pe_scale,
-             "e2m1",
-             fast_math=True,
-             acc=qk,
-         )
-         qk = qk * logits_scale
-         qk = tl.where(mask_h[:, None] & mask_n[None, :], qk, float("-inf"))
-  
-         n_e_max = tl.maximum(tl.max(qk, axis=1), e_max)
-         re_scale = tl.math.exp2(e_max - n_e_max)
-         p = tl.math.exp2(qk - n_e_max[:, None])
-  
-         acc *= re_scale[:, None]
-         latent_scale_t = _decode_e8m0(tl.trans(latent_scale).to(tl.int32))
-         v = _decode_mxfp4_tile_with_scale(
-             latent_packed,
-             latent_scale_t,
-             BLOCK_SIZE_CONST=BLOCK_SIZE_CONST,
-             CHUNK_DIM_CONST=KV_LORA_RANK_CONST,
-         ) / v_scale
-         v = v.to(q_pe.dtype)
-         acc += tl.dot(p.to(v.dtype), tl.trans(v))
-         e_sum = e_sum * re_scale + tl.sum(p, axis=1)
-         e_max = n_e_max
-  
-     out_ptrs = o_ptr + ((pid_b * NUM_HEADS_CONST + offs_h[:, None]) * V_DIM_CONST + offs_v[None, :])
-     tl.store(
-         out_ptrs,
-         (acc * v_scale) / e_sum[:, None],
-         mask=mask_h[:, None] & mask_v[None, :],
-     )
-  
-  
  def custom_kernel(data: input_t) -> output_t:
-     q, kv_data, _, kv_indptr, config = data
-  
-     if int(config["q_seq_len"]) != 1:
-         raise RuntimeError("custom_kernel only supports q_seq_len=1")
-     if int(config["num_heads"]) != NUM_HEADS:
-         raise RuntimeError(f"custom_kernel expects num_heads={NUM_HEADS}")
-     if abs(float(config["sm_scale"]) - SM_SCALE) > 1e-6:
-         raise RuntimeError(f"custom_kernel expects sm_scale={SM_SCALE}")
-  
+     q, kv_data, _, _, config = data
+     # TODO: bf16 kv on small shapes
+     if "fp8" not in kv_data:
+         raise RuntimeError("submission currently expects kv_data['fp8']")
+     kv_fp8, kv_scale = kv_data["fp8"]
      batch_size = int(config["batch_size"])
      kv_seq_len = int(config["kv_seq_len"])
-     q = q.contiguous().view(batch_size, NUM_HEADS, QK_DIM)
-     o = torch.empty((batch_size, NUM_HEADS, V_DIM), device=q.device, dtype=torch.bfloat16)
-  
-     num_warps = 4
-     num_stages = 3
-     waves_per_eu = 0
-     split_kv = 0
-     if kv_seq_len >= SPLITKV_KV:
-         if batch_size <= SPLITKV_BATCH_1:
-             split_kv = 8
-         elif batch_size <= SPLITKV_BATCH_2:
-             split_kv = 4
-         elif batch_size <= SPLITKV_BATCH_3:
-             split_kv = 4
-  
-     grid = (batch_size, triton.cdiv(NUM_HEADS, BLOCK_H))
-     if USE_MXFP4 and "mxfp4" in kv_data:
-         kv_packed, kv_scale = kv_data["mxfp4"]
-         _, kv_fp8_scale = kv_data["fp8"]
-         kv_packed = kv_packed.contiguous().view(-1, PACKED_QK_DIM).view(torch.uint8)
-         kv_scale = kv_scale[:, :MXFP4_BLOCKS].contiguous().view(-1, MXFP4_BLOCKS).view(torch.uint8)
-         _mla_decode_mxfp4_kernel[grid](
-             q,
-             kv_packed,
-             kv_scale,
-             kv_fp8_scale,
-             kv_indptr,
-             o,
-             NUM_HEADS_CONST=NUM_HEADS,
-             KV_LORA_RANK_CONST=KV_LORA_RANK,
-             QK_ROPE_HEAD_DIM_CONST=QK_ROPE_HEAD_DIM,
-             QK_DIM_CONST=QK_DIM,
-             V_DIM_CONST=V_DIM,
-             SM_SCALE_CONST=SM_SCALE,
-             MAX_KV_CONST=kv_seq_len,
-             BLOCK_H_CONST=BLOCK_H,
-             BLOCK_N_CONST=BLOCK_N,
-             BLOCK_C_CONST=BLOCK_C,
-             BLOCK_R_CONST=BLOCK_R,
-             PACKED_DIM_CONST=PACKED_QK_DIM,
-             BLOCKS_CONST=MXFP4_BLOCKS,
-             BLOCK_SIZE_CONST=MXFP4_BLOCK,
-             num_warps=num_warps,
-             num_stages=num_stages,
-             waves_per_eu=waves_per_eu,
-             matrix_instr_nonkdim=16,
-         )
-     else:
-         if "fp8" not in kv_data:
-             raise RuntimeError("custom_kernel expects kv_data['fp8'] or kv_data['mxfp4']")
-         kv, kv_scale = kv_data["fp8"]
-         kv = kv.contiguous().view(-1, QK_DIM)
-         if split_kv > 1:
-             partial_acc = torch.empty(
-                 (batch_size, split_kv, NUM_HEADS, V_DIM),
-                 device=q.device,
-                 dtype=torch.float32,
-             )
-             partial_max = torch.empty(
-                 (batch_size, split_kv, NUM_HEADS),
-                 device=q.device,
-                 dtype=torch.float32,
-             )
-             partial_sum = torch.empty(
-                 (batch_size, split_kv, NUM_HEADS),
-                 device=q.device,
-                 dtype=torch.float32,
-             )
-             split_grid = (batch_size, triton.cdiv(NUM_HEADS, BLOCK_H), split_kv)
-             max_split_kv = triton.cdiv(kv_seq_len, split_kv * BLOCK_N) * BLOCK_N
-             _mla_decode_splitkv_kernel[split_grid](
-                 q,
-                 kv,
-                 kv_indptr,
-                 partial_acc,
-                 partial_max,
-                 partial_sum,
-                 kv_scale,
-                 NUM_SPLITS_CONST=split_kv,
-                 NUM_HEADS_CONST=NUM_HEADS,
-                 KV_LORA_RANK_CONST=KV_LORA_RANK,
-                 QK_ROPE_HEAD_DIM_CONST=QK_ROPE_HEAD_DIM,
-                 QK_DIM_CONST=QK_DIM,
-                 V_DIM_CONST=V_DIM,
-                 SM_SCALE_CONST=SM_SCALE,
-                 MAX_SPLIT_KV_CONST=max_split_kv,
-                 BLOCK_H_CONST=BLOCK_H,
-                 BLOCK_N_CONST=BLOCK_N,
-                 BLOCK_C_CONST=BLOCK_C,
-                 BLOCK_R_CONST=BLOCK_R,
-                 num_warps=num_warps,
-                 num_stages=2,
-                 waves_per_eu=waves_per_eu,
-                 matrix_instr_nonkdim=16,
-             )
-             _mla_reduce_splitkv_kernel[grid](
-                 partial_acc,
-                 partial_max,
-                 partial_sum,
-                 o,
-                 kv_scale,
-                 NUM_SPLITS_CONST=split_kv,
-                 NUM_HEADS_CONST=NUM_HEADS,
-                 V_DIM_CONST=V_DIM,
-                 BLOCK_H_CONST=BLOCK_H,
-                 num_warps=4,
-                 num_stages=2,
-                 waves_per_eu=0,
-                 matrix_instr_nonkdim=16,
-             )
-         else:
-             _mla_decode_kernel[grid](
-                 q,
-                 kv,
-                 kv_indptr,
-                 o,
-                 kv_scale,
-                 NUM_HEADS_CONST=NUM_HEADS,
-                 KV_LORA_RANK_CONST=KV_LORA_RANK,
-                 QK_ROPE_HEAD_DIM_CONST=QK_ROPE_HEAD_DIM,
-                 QK_DIM_CONST=QK_DIM,
-                 V_DIM_CONST=V_DIM,
-                 SM_SCALE_CONST=SM_SCALE,
-                 MAX_KV_CONST=kv_seq_len,
-                 BLOCK_H_CONST=BLOCK_H,
-                 BLOCK_N_CONST=BLOCK_N,
-                 BLOCK_C_CONST=BLOCK_C,
-                 BLOCK_R_CONST=BLOCK_R,
-                 num_warps=num_warps,
-                 num_stages=num_stages,
-                 waves_per_eu=waves_per_eu,
-                 matrix_instr_nonkdim=16,
-             )
-     return o.contiguous()
+     splitkv = _choose_splitkv(batch_size, kv_seq_len)
+     final_output, _, _ = _module.mla_decode_h16_fp8_fp8(
+         q.contiguous(),
+         kv_fp8.contiguous(),
+         kv_scale.contiguous(),
+         splitkv,
+         float(config["sm_scale"]),
+         _hsaco_tensor if _hsaco_tensor.numel() else None,
+     )
+     return final_output.contiguous()