submission 665985

"""
HIP Module Dispatch v2: hipModuleLaunchKernel with per-call output allocation.

Runs a precompiled hsaco from a C++ dispatcher.
Pre-compiled to allow using local LLVM rather than older version on box.

Placeholders: KERNEL_EMBED, SHAPE_LIST
"""

from __future__ import annotations

import os
import sys
import tempfile

import torch
from torch import Tensor
from torch.utils.cpp_extension import load_inline

from task import input_t, output_t


# =========================================================================
# Pre-compiled code object
# =========================================================================

import base64 as _b64, lzma as _lz
_all = _lz.decompress(_b64.b64decode("/Td6WFoAAATm1rRGAgAhARwAAAAQz1jM4QjvLcZdAD+RRYRoPYmvuNNpsKdQSzUFBNubUH0irBGGzYsJpF5IQrF+doEpFnVsqZox2jHoEq6Z7CZfH25rPiZJ+8DevtRLFUMTQqtyoF+M4L4VBGPSsazlXrhBU6qg/5qQ9u1YsbuX7V8uW02M7jDIOQjljgyUOEpXTziMbRu4zLbXnnRaPh0SViEBLaFXGvHdUpBAYIh9NMznxqtF/kBVsoafLfhqq2Qu9vGllZB0ivbiBo/F2rcFXepf8TSkDfB2GLg4LQ5RXZ2TkY+PVymy8VK0h5+KlYWCXcswWNdp6j5dzGPFUaIqyZ6LNchv9LxYfpSlbeFiVo0QLpmQ/P/b/sElRKqK6AMDpE4o8m6Q3uPO/8djzOGzSvsC1yG6diZm8D6No4EjjdQm2c9X6WpjUD0lvF4DeNIJS35vbPDN3NMxBqC6mrIVW+DdQ3yuvWviDnXXUXSX2937avMMds5UrqxyODB1Yer5EVO6yoY/hnksvOOCqgHDXe3WM8+BrqS8Lp7T7O6+m5Z4fbDDbkoXn/4+ioIx/TnsHkW46KRXobyNv8pXo7rJ2ciU7Fa20EAG546OpJpZqg4ffInhIPCx01XcE5RGur8aVC9RV7dUHe6Ssn7pxdJaR3wy7ZRuECEZWpvCJIfgGGe1L+5ZhV/TpvP4Yh1gu2/V294Htp00wUjnQQ0KKJy2MXt1Af1tUYFrlbQaBgjN5zbu8zVuGbyC1PtF2yOfee4xNZsefNxFqVz5ul1/mcheBkeWjAPBlYkm/gsZGDsIx85pamCVIHygmBh5VVfnjgFhACbEiXS595asyn/wPDpIDVka/UO4t5KknHN1zDatH1Q4/QqoceoLpUvdaHSHQCcXIwzuXcGmgTtsKy1tXyibzQk/SqqNeQlIBQc6NYDTCVgi20M85D4AXHph2vUyrWavf5WisoyvEGnF9jNxDBJwPUPLpqhk8L6oojslvPDjcInNYmiTytD6/kZj6QoWZdyzACW2yHMHwkqnSHWnSAyvcxE6xjz7zxa9ExH4xLdDzRTsmsF0CLy4YvuQgK+Hb5HIua+wzo8Pnuugw0PjQDOsr59aj1Z/nTF6aJJDsu9EuPwT9P1cB/BdLQ7yfxHUMN/VDfmes5Z6ahaTBIJzicUjn35VyOfbPvGEMvtBDlDgex8PTfj8NCRWmCYKRyGkMKVJ7VLoUy6A85R8wT9bcGM54tQ5ylYdjKDE5RgzzqepkNh27TFQ91X+ehGejtsoB0YIEM4ECnar4IxmXF0+qNCm1CUvOhpo/Yb182RJPxp9XR20JJah3Jc7xR8pn+jAD2PUJQ0oMtx2fOO/YH5Y0VLnnZS6vWlYJQuCWHkgyvhH0BE+es+d2RWwuXItUck+XBBFFxBg7imFRB2y4dRVspaZkUtbn1r/e81+BNTKCavAMbFXTckYTS1FuvzFuioWAkiLMfMK1gdCUXNcSeOvFJF05BPxEmaGBW+N8MwXzl6dmQ+Cl0B6xkRAhftpb40Ye0G3GmlR9zEJ0V6xUW4lrznZPoSuYBTHH8FFu9K+2/nm9+zX0KawQjXVNalpU2oAaLvVIoENpBa61PqLQ6VDUNnqoWYuHh36OJL0HSj01bZzojosVew5gs5ZKwviL25a8B1ljPNI767SGFm77HClDf+beA20eafR094mAi4BiQP9Vdf0n1TCMBQXq30Kf0FAqQw+LSFV5W5+O6kYFtgAwVOVGPrs3OTjmiZlQaAglWVt3FlYM9Q7zOBsUBFgWl4QPSOjfoNNthKk6uMD4SK6DwiPQ0NqXlLieZ3vvLs7LtjMqtnnL6kLsD1sYzT+jotaZCzTLYL/OhaB5CcboPPZWjPH5Y8Kib4QGvE2CevxRL88DvFn+Y+b3Fl2Huit/LF6UDPVvB2dOo/ZiEIx+ijgVhxo53lCwDx5vOi7T9AOapkqYT39Ytq5l4MAJeHs70bFkwmNcbgZ39shFXgwAloSYWSXv9Cze38eQkGQi1g7NGKBDUG0+dMp1y5aCAsElLTAVg5sShr2ERIA/SvJOElA75srynBTE9dHhAHfU1U8ibL7KwJ8A2InQs+MOvTE6MiFshI6nQPndIv80Lu/FqhN3gp0LaHGhpaxLcqMz8kG3Wz9HTUAPouMsEJV3O5Z5JDjTquQD+Qpyeo91uFe9OqiMJGW4kd/t7n966xu5NQPLLE+ATj/LA+mcjNtCe7bKejR8ZiTw+aUJYXvWtswk/o/D/otBBig4xvjP/FRzSlOyJ7dSesSGwbkcRO7wwJBBZtMhHe2fSRvYTRIQjlO3Re5wyNpXCnLK5t0B1KYbTLsREFzHmhLQxBkSrDoF1uNsrFcOXXcw1eZuM/YbkW/iVEVkLrauou37U7BWQ6lz5L2/B0kKY8tF99tD77l85MvfioSbUxp13PJ9hrsJbJUOcXiBw+N9eJe6i6OTkrGMh3yLtKoYXmdfTImvBXyZqZRrfQUlGwXaIUkaeNHm3ZVuSm7UrqjGZbJQsS+jRosjxL+fvWr0ofBxrWVyRw84jTqbgKM+ttr5agSBqSQ6TEMc1gUMB+0kBi+NfFG6eSqlXhyG4O0dEDhy5JT8kc2DaQBvn2ltOixLu88GCHABujrqXL6QVmkM/OOEwpbf0PMOq7zj7qwVUaKQMT7DzZforBEh8apfYD337PUiJ5YpB7wpxxLt28kdHJG7gUxePIumZuhlENt1Auavdz1rHjhzAasHNHJDhQCqfOATQpZrXauO/ckou6OFrpzWh0fa2XXxgeO5yTEcbG1BSGPg/p1lowEyP9W6bfoEMR5+tCNzui3j79ZiZBT2PS9NzCe707wxN29lTLoJAWXQgdndhZNbUwLV7zd2xtks4S3oyS/u0DjsrkFctHT3keeQk67F+c+4X5V1e9ElvUQeDY5y8s/q9llBmAYJ5/s0LAabf/Rrpd9M67CsP+nXdISJirUaM6QHMjX+5liokZabvrJovvGKGqBPrN3Zbr5SD3qA9rOB8JVBM2B8JcAnH27z2ZQI0DQA6oLqgRt/nbF1P9aSXL3NalMdWDLjZIWDurVcvsMPt6zHgemeES8XX8W4VjoJie+tOeHyagv0no0psr5ezufTW45M5Y/nYRaCOZqHInrlbQc93gIaglACrfN/mjikx4QJKXb3DBB8I5rQ9Jg4RGYHQuh5eTZF42mrw6WcFUSAKF+1r8QEeryUYGmzlzPl+GQP8u0oeVnq/GfHUd4b7oENhe2a9FrmNvAjCIAjvyTJbpLwCP0sDTBrQ40c6RoBXhpBwimBEO0xkDtPDuRK4/G97VKoVbBBFfa7Uor6vUQWkdhB9vsNe0lbFX9MG4KRIH6aKD2xDkKnsWCRX+2trGWaYW0rmVt0/GbGKk+pGahXQlLP66oTJP28L7KPRCJPtvztrupUQILawxgQrRPjMY05gPeza6uOOIAcrX9qhvBdG3e+W9hX0oOtTKY8pBCFvYjFM9XqEncMyRCWjy2xalXkVPOsISB02vaOs5pGgeW1YXVaj6Ehp9pt+QuTIucpeSK16X8+QhPkxdznuxDz4UD88pG1wgkmrwoQCYG1uMlE29HzmlQ4cBudtrN0Xyxc0K8OQJ/f5SQBtEr/HnzKhHp+nE73QCYnqkiO9OFfV7JqJlHtL7Mml//RpMTjh4993ui7OmgFJrg0uJEtCpuW2ufemVxlEhNH2r8iAKwVhhAWtRHsxXBZTqqM8L06i2Zg5SucdYUyn+2MNKovOfpBIsh5LwX3eUr89+FaUtKktOMQO4Y0MtCq8aB3vMBRsQwwbgTobSODu5GI0LLWkIhAyZV8atjrSImSr2oUulMvjTEgoX6PJ7SbT1Qbll3XO7nA2mzV535dTLtDlxkYMvzLI+IifZfL6o2uuBaFdEzbfiCKnDQ7LukWLTJbtJBTE6y0jzXYKNee2Ojrha1Xt6SEDI0sllnERWHBuajhjM//GUH54pYXrR4rNHsoTDxgMIzyQaX6xvwWv+k6nYhorB5xmTiUiam22MUIQICLbhKVgyK4nDRetBOkgQ+BjOkn2gJn6qmNMbv+ri66KyWXJVAh17vXkUzpfvLi3rszeBwnEIbz0/rYosr+hp0IvtWCRP4mfcuQjvyIW+dB9AIRv6+3pWAy1QZbka/Cl8tzZ4D/rz2+FYHTGV6KywIv+9W9LDKabUvouBL+KLhY+qjonSj43oISjRjbJyocQJDcmRQCGnza7u6BKlgkaNuRtYIDoKTU3tGFuBY76/9TgimZqNJUWoEtVLVtCfNU22YL+UUW+sasm1/aBCdEdnvLxoPzFR6cWSYoixlDFkuOhmwD+vDP+nOdWRtshc8MAOlk1txsZ0GE6zyFgcMr2zrOVZK7aWMGch8gNJtTBAzIwdmTI1WbA9WHaJKurH6aQaTEEOquCZBMfFcSyxYzrAFQ8uwxJygTbgFeBRN0IRyq9PqoGdMo+a3h3ysnmubFD4CkVn0UbKAI3i+SFdzji3MTYjFfwCuCzseYxmzSKUzbaYIBiAgZob2O4n8moE/d0qz+FoPBu7FwUWejNQkaj2ug2Fp+b/xozg2/AdQ8iFHtGlotnB00R69seMmGgBadinxD+GePpJ9KWH0kCi7jstIWiIdB8Vls2bOt9t+B6AgrmM1nCQntese6AZa1Ksro+P2TTf3q1vUXmw3Il+cGQFczuKuftd7QqI+/AK/tU3S5z6ed66C0ZdyjhtgBw3NeZOygieHrmw6UG86Z0Pj+qrokfD1zAcQRrgNp/SxhfQWYkGG/hjgsE2nGB70l0wHWp5/x1YS/PiNF7QaMZzARNGiiVkuu6fNjXNaAIdm7WWT++soj8MwWfbVSHEb4Q8xKEY4OD0zCeBydiqVYffAAhYZLrDO6YXiVdCrP1l/otMq8IMIC5hBWpwmjJsPnANojpRXBsGSDf//AOiihrSmz7VuOrQoZextuNhtJmm2ldR4Gy0a8w1I9N/+5cjlogojfHdzbO0rvXQqswvwwUIkqoPBRpaSJd9QDVVwjXjxfQzIrKUg5axr/j3V9RFQh1WQSYdrNqVKHo9KlKgU2NdaIwaR3dmWDobfCbd5t4iaqNM43EYBDi7ViiIYxYWAyp/Vrf05ik/M4DbAW9elz+N70L3Zy04bAYko2qD/+0n+Ef6FPqTWdDilZSPiWrXKC5k76jyJY//1a+Bu/JccD0IVYCFjwiwLHAGd9uuxeudQWaM0oxROdrRR0qk+U07rHDRqDOb+11IVZElWKSqs4IJsteYex6quDDGJof2aOB4NDqUOGc09xF25nhK2JS9DLVarMqxrDik1sOHU6vj0QpPCtgD3jj7FMGd/xFor6qyxP3gpmd5NCQXyH9Sq7UPtPoBPuS7uXT5hVTUFKUMPTNSctOEvCoOUcIOPtNBhG/fdZj+vq7wNPYry021bBJhOWzXwPR4bDLlvu+inn76BxNSH5NvaOeeK6eBZn1CvEnm0gL8GP58ITWzaLJnuWzRGeK4caSGffskNxYMSxSNvKtJa5srSuu/fDf4FyNvufOqJuhUHP57qtkdGYel0g9axr1iRbUyBxBFRvgC4c383Ra+plBGZa87+tMkdK0ddrC09AmbVgdhzhMB/wOMI8lOmVnguhlwI2nmTv147RnM+Z77SNwkGRYJqXEH44OODbRpNe6UUBEm5iiR3BdQM1N/XSh6exbT+zpSCMlokG8vUSaWxpnJhUwSfk77CUQHZ+89aSv1XLNQTy1v8x/WX3fv7+wrU9leNQaPmb50Goxkl5RJaDo9Ec7f6RF+BUB0vbj/VvECsyZNzac20EL2PyyNC80w0dYC1Qp1t1x2NGp6Vj0qIK34bRb805x/gAoaRubjqqHqc8ND5WPg3qjJ12/I8HOBMueCvPO7s2fDv91z1+47iWfJaj/HM9KMSI51vrlYUU4sGrxV+0mnpWMTiIWv/6tK7pAkKSCLeMUEyjIfNi1G7jimqIRtMvjWf1k8IjBWQA3tTaBD2Kg1LIWWt+K7teQ48/MOCn5p+Nfcq3bS7yjk9eQnu4roHn3Zh8Tb25i1NiI7PpLert+xlH5s4V9MJHYcg6Y9SbdRauo6jr7fYO9JBdQNBj3+36vgZtNwMLnna+CgKFXg2MBAb89yjpGxUNQ4s8QiEJ/4vynpJm3XawcCw+l0cjur+j1FKpTwASYJSVu2/Dh3YJzCFSIm6avQ+q/LLY3g6+xgaOsk9eC+wajN+rwYfvf90EKVeTiyEr8MJ6LEjK7tL/njLrXBSYsUFilJPqafh92/OQrc/+VNIWOxibx4mWOv3BS6Sbp/xEUydGXHetySKCnZGMKPDiC7P/XVsDG+MtwwlVoGIVtlz4pCRdYqrdjm/pbEB3B0sGtEzKpXM+34KDL19xEQmecUzWsrrKkn/Pz84Spc6wXuW7k/MPaYDrYPwMI2/v/KFMhmWiIjFQJz3YcIzBJ5TyMqlNGg5yC2Ku8IAJDijJmr85nmzv0d91xDeGEyqkc7zFBtE9Tkd1CEXGw+hu+L2S6B0m0zjVZ9eIOM5wV/04pjqiwDH8OPAxWWyf1c9gb9YHYEZY1X0dynbicIIvSHxFjB2lMHxnOS7WbrsIKWvDAlKmUeSxRMPnhV+kvjv7Ic2mMVjNEtavNJH3Pe5AYI/LFlfBdzhM+FLd71kokdiJLyR/2yAoGuH7pF2hvJSq9CtgLvs4u1i9vq8H8OilmHCzvMqfRTN8lDFEi7Ahmx2J9Dl81Y8e6vMaezfWiGg9y6Y2LuxAv1u8MA5Ggz+1BAx0654QYdhmHUyFf6yYHFFMNoN2Db9VoRdmTA7Jb6L+O+qfGUeby7ra8Np/rKuoZixffckcmY2WcU3Z5zKULeqV7HjNx3NbSC+M3yWh9vEwQdSfSWz2JnF5DskxneFmPaLIVsyaJlpm67ABJFFiDFrK9cejtsjhWkSHGn0zwAPNckeINom1TdVD715aqBv0k7GQDRmlc32eq1qztgjSnpE3fr4fgQbUbrk0M3HGzcdfmOxDtZ9QPLpt1BExgwRyu8qwHuRP/7G5DujJqJUNEJdaocdE3NIFQVGJRWv78WXYyGW1cFSJoWDkAeWCDzGZN1OFq37wzAI0PV3w7RkIqV8RJ3w/+9VfmdZmbjmBGITb6i7l/MaDfF1KdeM9Z0qII95yT7TbRb3kvTik9wUsZbSMS9/7V2whB1eSM2JBvxqytmKFp4SaVZP7zSoAOa+mu3bpUiySzV+KKaBNlBAkjG4wnlJB9bUfKfFtYnHYnNO/4YnkDAXgBgj4JF9M51rh6e13WSqWWTacPFMlPH1WZbEJoeyK5IKPXWzcPxpscame2liNCaXisVg5raIaedqnuV9ZN4tqNeG9z5Ew4dzU8ud0vUDjE3+h5asB2PPIq4TyjJPjsbiNyKhtAXeL9ZFbZzWUPRUzJVeEsqcx+j7EqIOb1GxsCeXIG/JBU1DdvLOYTHafkIFUEPT2u8WB61r1TohJtDmkJlQkQfZ+FXk3QlAlrJfTNCdtQVaK8zWbGHkMQfNt+4faP2moCHqEeK4bqiARFBovAymN/OYtxAvqwWCAO8AO6K30bY/EZN74HE+Yi+7i3bKzbAYH4SJyTB8gPgWzGp0NFX8Ej3GBbQ4SKbPw6QZHTa1DBErLFuUls1fxjgzrYnDcav2CeEM2EtUL9m+36q71JZU6SOZfbmvk+jzAvqBd3myT5ogHgmhJlQvoF/gRxTuR9cqI3nD7ihXTtn1t+dv9L1aQllgj10M90qRmIGZrn1qaEXwWDEIxfnLfRAPENWLJUg+SPRoPA+gaXWCIxNFVUeDFZIZr5xpkEDUKNFPGWN9q0YzKdD0Tvjr3ViCbDpNi79XnYRjhtiilfKl4mB+AoEz4MtR/xUzFm1O9Wm/hYTEfagxG1rDqvwnmAfyYjuxNP6hSXz5QcWtiIrVc8xWnUf9Y4HbWQtLUMwvu6J1JVV7PX3hreVWNj3KSB/c1l9P0keZHJeqz8ofmzJedQoL6dVaFJIVe7nFO/oYpDBiUCXKeol94PJwxA8BQeV2caiGEofE8DXnMJPbpzexTmfSq/klTUTBQz9hpiob3sBY+kQe7N3AhcCnjT1/TqOw4jh4A4pVuPZRw86QIYtRg7+2oK73NxN3o50GcbtEMjAnb3Q3PWptC5XzZn9G1atvI5KcuLCokfP3w24DQPflLSXdXcqhjKUqfahf/fiQMcd21G48D8gycCT9nYPSGhsQApVLsrFLiOvJJg2y6Tr76AQw9VWh2a7BPWq9ykyasvaDP1OES01HqxVsPQbcAvo+56DvvRU8UzE9Kjlx0Lr3AkVVmAmCP2gNQ4mp5F8+XMTWjJmONhiPwLiOuztAiqcu97gYzo0Q4WNIycaRUrd6T0927wAqPpnfoZc/5tSTSLIJL4CmTlHyWc6Dh2wZxB4w5X0lZ+PwQIBsIUi89V76NqRfmyS8VDL5RiOGzDx/Nl5mvNgk+imIi2zI4x41M8/GohBZkfcwWxkSYGOKysCtmOUhA2iMlcUNTfkVouEDjedwZUVr95hMCaquTfqcltS2LcMdgLOqsWb+GxPHq2wMKojnrgjST+8PB9jR8NiIL7cqHDSFZa+z+KReQYMDaUr7C5NZqkyRaimA5CXDd7XIsNS3rn4zd3TkJPI+dW2bTFVaFl6a4ZqOyNqu46sbIEWhWTGI9YGUUwMRwJbn56r7KLV8tbL+Tk+YlBhhAddYmI3A2wqGi/qDkK1hCePTQOoXzC+k6XBoYv9XhrcD0Ue2AxsmKNS9zooe+glF6U3uo5psMoFc1jpbFhXC9ow/3bEVz1huyNh18a5g1e+eDWbsyj+l01EPlzwYwkRgrjhz+to+W9SLzU41TJYrOkbKE64y3Sp9d05ix2p0mvlzcTeyu792CI5qMc7b/u+YDJ3AkMZcXy6pzp9VcX83DVRaR2xYunfC1EYa8CNlvKxejQ6v3n9LS52i6KVKOWvJcHWji9tuVu3Px2uqNUll1Zy6mU3q77WamIrkngk8JZsFR2cByOK8wZ+apKENXU7eW0YKSaivnwkjhE+j8YeUBqIq5PV3Pmq3NidZzOJVip4vxh/nGFTwHSnv8gI2dLbq+G2ndk4pKLN+C0l9xeGcsQ1f/dtyQWrodBPbtCKCKdEEFTE/v+fw8bfIVbBEmzPgk6QBq7r2RqMyvoWQNFmRA5iFiTPVESCgs8uYgW8T7F73ebENPqPcGWTYANp7rljS4KG2mTm8k/r0J4eviwbCcQwRp4M8tkGOhPb1Z9qJBdWH5fwQobU6zFeXCV8Nrb3s8R7nfo1WvboK9MPi60L8vd+uC32wqJ37hRFyZqseFkCvjUdmISGCwBeWJDETzltg3HwJc9Bm6I39TE0QfOSTb1AwSdWB2a9xmjqP64WzXq8SDe24GHMHJzHmsQVOGDjf///j+ukD99u2aT1fDygXYyi1dMy14BahHW8a2RAxHDCimo86ZC0u1iCYivjUcplYKgAmySQxTMYD77IpkrVCFPp5Lnl3QwqNZ3/5YAR56gZ00c4/+RL37MbWCSgI7HESEg/8kfY0h6Va6qSrCyAJ8xLsssCZ3N2xoJHwFb7pbXRpDQDyEOhbF+nANIQlwWzKwyIa+4Tk4WxA4LIypNTKOszxDIBATQXvHTcOWkquOX6gedo456TR0KCJc2NvYdhvLzLPahgVJhZuaZVc9/++7qxTPB5ySolXzn0VgtkEYY2BCQs89DpGoxkgIzl+0PyNXNb3C4UaV8mqymOWAUla+d3jgy/ceqY13cfq2q5FaH8NjKKPDEUEvGraRASPUeZe/YFF94WiLAafT4Zk8KLQEPJpVLFvSqIn7QF7bJ/Uv/ma4g8eD3bSZiSKsXDOiUEmy/4oCSRfg1oyNfFR7DNPpUQR9H0VGdaD3ieSLlGdQ3aTYazF+5E3kw+Ynh+5B8AXsFAt7ezD35HNQHIuV7MgoH3njQVumPJCojdp8MLNvH1Ws/tsy/Ji1DHEZprtmqqRyfjsGI4X/o0niJ4cD/WONw1INCd181dDW4J5b1OheXlZW7cFcd8kf/p02RL52/pcK8hQrJZcTXneVLwJYeZgbR11SpHpH9WwVtzGbHf8YyuUN9gM7axxynVEFvLNp+OMi6TVLP0uJ0evZUERDE7dGo7nlDSvqAUiZe9Ryh8aX9Inm6qKkbaLmnqw1enTa4QXK6OEcWyn3G6SNqYJrcWtPslh1W6xhSHtgQJh4xf4sszaoaBDlTmCQNQKKctnQL/Tu4wdHCYaRiHvRG5UHFLQANDQpdErlHuSK/iSn3GihQllUI4edUpKurMUjfAQAHKukvrOUlEZFyZWBBBKO5fb4xKY6JbYlagFhD7Fe6T9gmC13fj05dMnjgdMwSI+kGAIXm10jhoU6KwWkPdnwEvjnLlWwqT6mdhhlmHnFNHl9A/MQPgypcRL/MPMhyX+Jnr6NqSmDSCZP0eb5APjqZT7ORASUWPGc2k7q3TRoGKcYuXwyeh4ZKh921q06LSOKS0t37l/hAAEQA4rpxCumPupibpsJ0qcJ/yHXDahzlnuCoOmGUZeD0fS+pbN03w96w7bZd5LDZxtGlVzs1VPt2+Se50J09z0SRno5xG/jL7vIuXZcKOfrd5U6+NieLCb53vr9xMiYKsK8/KsxIZXZiG6AF1/aH6p45SMKMZlPz4ygZ/1NvdsaDzN0juePnpUmHDrwUs3vUJQWx56zBQrUYBraUyhctxwIE0qvlwJ0u39pBT445+tQ1aRHXAdD2E+qjfUuguGA+H4ObyUGfR9gwYy10hftcwbMlDlQZdFqFq8lhgv6lJDiORhR6hfNKqtuJXjZxZyS7ZG5xUhT/FdeDc2/Q7UsFhv6YNYGiDx1rIZ9d0HtYGSdWOULoFG28vedtcYtRO3VQQ0NR1+m2O0C2yr+upYFmUz1aubFbMcxqZEcUmZcUAofne60juzs2SW647fiM93YqxBMiXVxd09P5LnHKNa1HMBaTRWfTeAvCsEqYbW+OBIutZS686YTCXaP+9RpREUJXsZzt2Fmwt9sU7z3seCeYtAhInf/zKRp4xm5EHpUE6eO4PQUWZ+uYEVXtQro0WlQqRr1c5ZnheljqOG79uQMs1rEA1Rxr1BYGvFL9Uf2EDUECRnLtTQq8+Zv7Dmg5iAgNq59swiEnurBmqiLUIAWP4SPyRrknpysC1MAuaTsGl4xPC8qSq5X44LBdplZEKtIwubTlocsXLC0CIs+H8Bk1SgMNPC2UqvqWLSuJcJdw1eG9myPxNRHC6gp/WET+FK47PLyPDJ7Zah9/lxj6mN1J5niYTacbeQVMgYnTtkS4l35bKDiWDEZ+6pc1ZZzMPH5h1j0tedm0OeX2/2sDMjaYNTuaduIcRnVRGA7zzDm+LIq5YJSTAaHDKOJrm8oBXbjiibQ5CdpxjCV7f17RrsW11mDVWsurhxkeOCUvtzSRTiYWnp0dlxvACvzhJlq0eiljo0oodYz5yfLeHmD2HB+W/OvFbrwWGgwbuFw3RTDqgx06zLuwzIDzWRLlS9VUEyTu5Ta1+qJkcet4SofLLXeMR6zHS8+i+GRXau+sVsjcfYktba8cn6UI+h6zLm31gScyAktBnJrVpTncaOqFVXxJ8Z9pFiOGMOIwz/+fslWDu5q5nc+ezsAqySFA5XFX+byt9qYcjrvT5Ye4PhExwL70VyY7fkuwJUfhcFkkL97YgPfr90xmaQ0Z5WjHGFFnJBcXWoM3wPNZlSY/E4RTMJdQuzHojJJk8QIjXSTzkHJl7pxy0jWaqKndw6K/CXD05AYqJ3M5cYVhf35G009e1fIa+AuVR80qRtWKdjYC3C//uJWyquJgqpdeHV9EKbi1YLwZz1EQwVQVXnVrsI536oLdRrHo5vCEfRv8B4CCHE3O1oyE08dciEFIV57j4fVorI2xCS3pbDqaXYtq7Euk6EIc15OR35V3jbOYTHBV4V+bH+VbVF6vq/YGSARIBg1HxuGdkCUWxjvu40OoHbcFF9UHLHhjzk/1WkEZGdzC2JvjVeBfBlEKnunDCDy29h1y2L4lYLtTZmSamn0eLZ5dP/+Y4ZMrCDDruKQgSy3etYw3hZ286SpKHG0yPCrEsX6P+rL+PE9Ajlu9xn827OgLTBRRrLRze8vKqXfymhUK06HO+12DMBYxUO6+vxG0Rj0iKhevfEC4YypHbvmKhPMfCTTO+yqu4HjX/WLTrZvKgwTnGl/alMsj431VTr6hKk0PRtWAzlcrZdWsJ0e6zJztHWl+hJdWnmOaZgtE6u0KmAvPXzQ75kW7qQJQHjP64KixH5anWyYHXPm+vhP+hjf6CWU9GlpRQkZEPonBXQgbOQCdxUYcSXSAVLK+ab52f0ra4GRRHeJj8HxmvmGwS7Ao328dOAFwUTp/Ocz3vYwQFq6onLMdkizT9dtJ5rABxip+UUaZdg9OZeilCWSEsKPsIeTZdO6F5sCM038gqLHWIb/mSyqjz6u2OHn9Q1dejgWrEepo3twSzob7WxHtG4htJ6fnVMEYtqKGuQ8/+pW4jcp61G+cqWkJAatKpwylQTfpP3W0+uvYsytbwkUkgHOcbteEatfHhpRAvqaFc1yykZY3H/fXy2tjD4mqGlv9iSHwHGQrr6br3w+TU4czMqUVv1agYeKqZcXzVPKlOvaPHztYlARTpG87yqly4TOISqWk5Cmq91y1KRfWZSznw5kwq8hTg3WSWIxPgthD0GCrsLdVB8YPutyas4ONwXsDk8wtdfR66NmLBGaPDvt/QtaDUZuXHCnTdPjjZ/TpCs+oNHOErhoUCOd87S8HW0Wb+txrVS+Ie8IYzdUJU2ytsD4GnZMPFuqG3IlR2H39j+Kljsp7H2bWJjO5py9Lz3FkT00o9AUEqqEruSEfrx+qwtkfnENVQEjmDGgjQIPV+JKirZE4R1CDq+aRCBYc71upvAbS2uij2uh9Zy8kVM4myoju3IlkMpKN44thnVFFw9r8t3lM63CYKO0mXCElHnJA84W+Ch1T+JJuffduTOe8K4yTqAmPSVnotUt8eH+bVFbleZJzGRYEU2dK9rrRtmJ/PrCHWyNBVS0ueYdghGOk5zdkQbozM/8LfHnjbI8haBg9qO+h4dpuK+FWNKSmsEHw7sNV0UXrcoXE7fE4vWQBUoVUmz5PFf9CsTJXbx6J7RMcHPiqwsEMXbE+BbqWcmYHC2W7ewukKLcAAhHotMGFL8cRMNX6d/wWyMORS4cg4I4Dy1qJnPjUhYA73xQt8ABsbUMiSN2AxJ6WAr96myzHNAj3XuH7ZpusHRFE9BXMYjVO3vW1WdzGGpZxlD9qRgVCIAHlAasCxHNEhbG1bg952Y+2BAA2sMMIWiufwecaMSidtkjwtcOvT9hj8964N/UJ1Hxkf3EayrU15XgLmYJUXFcBTGj3PUkY5li9+ThQT2HAFHztRESJj72Z5fwS8g+fIs7Vb1yEqgtBrmzn0Jo6U9aevUoiBcdKiD4HvFHXbC7DcexHsqTl7Xn5AEi+yGEpluLEAfLIrLyWwROr+pJvRHTOqJUlUYkY4saVwFkgk1VQB7qxhS+ZFh2cpCokXFQL9AepG346di/19UyyW/m/a5F0QpIKy/O9Uxy9lprReECeNzBCyX4KRYYiHJtIFIFnProXxlSjm+7el2o7g8pkrJEbvEkreyR/bDuqPtX+UFcssg+ozVWverX+7eaSI7UT4Jx+k7a3qfm9Fa/lFbsjy7G3XtX6nJMkk0H2zxUaZGtnJSlKvE+V1209ZS64cu9tzSgv/FokEAtJsM0+Cb9fI+U0HlX7VJrnODxxagK11a+PJ8xFFHkSOIsypp10GVzPhMUsKVI3oi/tS+GMkzyEb5+7LmtkZ+DOOyPNqnbhLFHHO26K332TLLvLAYkMsYKuHIKSwUqIUqH/3H8dv/h1EbN9T15xh13TorgRoRpdh2VY7fSg1F/jrDr5a+syE7jQgQ75m0Z3l58NJqlHrjwqUlFb+rvOEvR9wsevmm1VUhJJIIcyXHtoUFt2V5nZ2E9Q/HUl4FtYx8hmi18npsy+6BvLjbTzKbWmvFjtAq8gqDNejXIoROd/6FEwv6IydF4sCsRPanbzdFjj1fn58zPRAG3tD+PXwNSK9Tl/JPbHR+sTV7bQ4GrJ/0PGjEQiygSt+2BiWLTbL+XvM6bJhPUrvMnV9UU8+jp4uijTodOG5vYg5tPDKT3czs8c/FmTa93oYVrzGQKe9s0+dIS+NflDNSl2Ft/lTxsYwZbX6BpeUlc5OnEyKOyCiL0Cscb6WA9oVVHiqM9veRCvq52pSPtmRthlXw/gKVD/vhAWUZEHVVORt3wbqSXcVfuNyULP04Z1Cvmos0cwnE4+s31Gf9o5IOvd7tj+Gx2iW2cy5t0mdLRw7G5oAI2iG64bhAM3Qoqr0+OR7JfOErIFJOSPUDoV4H2Ml3CUEZvj9IrYC8BhiF44DiFAOc3MU3XIbd9+Yo6fBpOb9XR73erxaN+VPLX+zG+eL1zsj3T6OUYPo8y0fT46eIEt8FpUa4mlZ+XdQCDUj7bK2Zqtu0syJO6lMDZoafaHAJp4JFrOiX9BLTysKaA8TpIOi3kqEuMVinzKpW16pJFmRZVwUyN4cz2mc3Ookz0MPV/jv1RYLhl+K0jMxPgaOuYftoCCsaX+LMzmioq46+cDnMV7jV0s21YUY1UxgQrQHM8KH+TdkVnAk5k2KQfBrdOQ1IA/l91DPaD26ew6L+t+bsW8cqf0RAJv57n5ZyCEYdEjnqKwGXhCrpU0yPH15dHZ3MoBNc3Wwq95Nkzb8tIA4npMhybQyHzjGAJrlBfitqa6s2YexLFWVA5sDSouQY0Ik137Flvel6eQa/FjrwJwUr5o9cdDjVTMuB2p+6+mTI68xWqwyFGwZt9dO11NOoMgZEAXonutoW6Bphgl81hVcBK0+seA4xGtctw7db9J1aoAfma0Xg04fG6guFTaVYp5Z6Wfr5zPXmnMMN9/wp3ICX4UU/RhRhjTpRn64sErjZ2j/OH+zH6FYwcKVY3GD+MnPbPndt2kiPie4c38erzeAefJe3xpGd4L5D+Dj06N2ZqeHl1jbrXP3ef9JiX/RbIQjvV755WbKdgXNZvsDuUojgPqQVAUA5DRWC4Y2/7cWEYh7ucv9xrgm/pOOAVfcmO6SMZheyYtvsRC4DUcwGS+ct1Wh2Ubkq5LH86Md4b5APm5iFfIDiPreiqcA6HzZNSKqmfrIwwrK0kkihERilhUia2T39FdWHrro6utsmFj1lOfr95t6eU3qc2YxubOKE95qGOeW/8uW0mLQsGtDfs05AKIst22LEDQ0azZ4G0blzaF8Rs5iBuxjtile3nFdvHdKTRy6RTz8TbytBW/8eozzMz5D2cxIg6F2OBaIbXluDR2OlAwxRh/o2M8Wr3oht4SyUMs20F78a3ZCV/PdAQMSUs7n/EZbC4CejucjT1GOj4K/OVcoTfB1vOPmbkfhELUoxF1UzJWHYmCkPHfXCi22ilOQOZX/NkNesw+XBDwtnv9Qva4mcVLdqRJbXS3eRcxPsUwYZyPvVMECfdHw/6KANCds12yCaFxDlnfaD2ui1eTKNzPZadKomeXUFIc6n1vZHbiP60LjqfAYLHhT4UrFafEZSgrnjY9NVhVcFNCE0nU9lAbJfYNWJWKj2hAAAAPeUxpn8K4E5AAHiW/CRBADqkcHJscRn+wIAAAAABFla"))  # 67824 bytes, xz 11784
_sizes = [8552, 12992, 11984, 10184, 10160, 13952]
_off = 0
_hsaco_bytes = _all[_off:_off+_sizes[0]]; _off += _sizes[0]
_hsaco_bytes_1 = _all[_off:_off+_sizes[1]]; _off += _sizes[1]
_hsaco_bytes_2 = _all[_off:_off+_sizes[2]]; _off += _sizes[2]
_hsaco_bytes_3 = _all[_off:_off+_sizes[3]]; _off += _sizes[3]
_hsaco_bytes_4 = _all[_off:_off+_sizes[4]]; _off += _sizes[4]
_hsaco_bytes_5 = _all[_off:_off+_sizes[5]]; _off += _sizes[5]

# =========================================================================
# HIP Module Dispatch C++ source
# =========================================================================

HIP_DISPATCH_SRC = r"""
#include <torch/extension.h>
#include <torch/csrc/autograd/python_variable.h>
#include <pybind11/pybind11.h>
#include <pybind11/stl.h>
namespace py = pybind11;
#include <hip/hip_runtime.h>
#include <cstdio>
#include <cstring>
#include <string>

#define HIP_CHECK(x) do { \
    hipError_t _e = (x); \
    if (_e != hipSuccess) { \
        fprintf(stderr, "HIP error %d (%s) at %s:%d\n", \
                (int)_e, hipGetErrorString(_e), __FILE__, __LINE__); \
        abort(); \
    } \
} while(0)

// ============================================================
// Shape entry — per-shape state cached in C++
// ============================================================

static constexpr int MAX_SHAPES = 32;

struct ShapeEntry {
    uint64_t key;
    hipFunction_t func;
    uint32_t grid_x;
    uint32_t block_x;
    uint32_t shared_mem;
    int m;
    int padded_m;
    int n;
};

static hipModule_t g_module = nullptr;
static ShapeEntry shapes[MAX_SHAPES];
static int num_shapes = 0;

// Fallback: stored Python callable for unsupported shapes
static py::function g_fallback_fn;
static bool g_has_fallback = false;

// ============================================================
// load_and_register: single-call init from Python
//   hsaco_bytes: raw .hsaco code object
//   shape_list:  [(m, n, k, "symbol", grid_x, block_x, padded_m, shared_mem), ...]
// ============================================================

void hip_load_and_register(py::bytes hsaco_bytes,
                           std::vector<std::tuple<int, int, int, std::string, int, int, int, int>> shape_list) {
    // Write hsaco to temp file (hipModuleLoad needs a path)
    std::string data = hsaco_bytes;
    char tmppath[] = "/tmp/hip_mod_hsaco_XXXXXX";
    int fd = mkstemp(tmppath);
    if (fd < 0) throw std::runtime_error("[hip_mod_v2] mkstemp failed");
    write(fd, data.data(), data.size());
    close(fd);

    hipError_t err = hipModuleLoad(&g_module, tmppath);
    unlink(tmppath);
    if (err != hipSuccess) {
        throw std::runtime_error(
            std::string("[hip_mod_v2] hipModuleLoad failed: ") + hipGetErrorString(err));
    }
    fprintf(stderr, "[hip_mod_v2] HIP module loaded (validation OK)\n");

    for (auto& [m, n, k, symbol, grid_x, block_x, padded_m, shared_mem] : shape_list) {
        if (num_shapes >= MAX_SHAPES) {
            throw std::runtime_error("[hip_mod_v2] Too many shapes (max " +
                                     std::to_string(MAX_SHAPES) + ")");
        }

        ShapeEntry* e = &shapes[num_shapes];
        err = hipModuleGetFunction(&e->func, g_module, symbol.c_str());
        if (err != hipSuccess) {
            throw std::runtime_error(
                "[hip_mod_v2] symbol lookup failed for shape " +
                std::to_string(m) + "x" + std::to_string(n) + "x" + std::to_string(k) +
                ": " + hipGetErrorString(err));
        }

        e->key = (uint64_t)m | ((uint64_t)n << 16) | ((uint64_t)k << 32);
        e->grid_x = grid_x;
        e->block_x = block_x;
        e->shared_mem = shared_mem;
        e->m = m;
        e->padded_m = padded_m;
        e->n = n;

        fprintf(stderr, "[hip_mod_v2] Shape %dx%dx%d: grid=%d block=%d padded_m=%d shared_mem=%d\n",
                m, n, k, grid_x, block_x, padded_m, shared_mem);
        num_shapes++;
    }

    fprintf(stderr, "[hip_mod_v2] Loaded %zu shapes. Ready.\n", shape_list.size());
}

// ============================================================
// Hot path: shape lookup + allocate output + hipModuleLaunchKernel
// ============================================================

static inline ShapeEntry* find_shape(uint64_t key) {
    for (int i = 0; i < num_shapes; i++)
        if (shapes[i].key == key) return &shapes[i];
    return nullptr;
}

// THPVariable_Unpack: borrowed reference to the underlying at::Tensor.
// No refcount bump, no heap alloc — just a pointer chase.
torch::Tensor hip_dispatch(py::tuple data) {
    const auto& A = THPVariable_Unpack(data[0].ptr());
    const auto& B = THPVariable_Unpack(data[1].ptr());
    // data[2] = B_q (unused)
    const auto& B_shuffle = THPVariable_Unpack(data[3].ptr());
    const auto& B_scale_sh = THPVariable_Unpack(data[4].ptr());

    int64_t m = A.size(0);
    int64_t k = A.size(1);
    int64_t n = B.size(0);
    uint64_t sk = (uint64_t)m | ((uint64_t)n << 16) | ((uint64_t)k << 32);

    auto* e = find_shape(sk);
    if (__builtin_expect(!e, 0)) {
        if (g_has_fallback)
            return g_fallback_fn(data).cast<torch::Tensor>();
        throw std::runtime_error(
            "[hip_mod_v2] unsupported shape " + std::to_string(m) + "x" +
            std::to_string(n) + "x" + std::to_string(k));
    }

    // Allocate fresh output each call (padded if needed)
    auto opts = torch::TensorOptions().dtype(torch::kBFloat16).device(torch::kCUDA);
    auto output = torch::empty({e->padded_m, e->n}, opts);

    const void* a = A.data_ptr();
    const void* bs = B_shuffle.data_ptr();
    const void* bsc = B_scale_sh.data_ptr();
    void* c = output.data_ptr();
    void* args[] = {&a, &bs, &bsc, &c};

    HIP_CHECK(hipModuleLaunchKernel(
        e->func,
        e->grid_x, 1, 1,
        e->block_x, 1, 1,
        e->shared_mem, nullptr,
        args, nullptr
    ));

    if (e->padded_m != e->m)
        return output.slice(0, 0, e->m);
    return output;
}

void hip_set_fallback(py::function fn) {
    g_fallback_fn = std::move(fn);
    g_has_fallback = true;
    fprintf(stderr, "[hip_mod_v2] Fallback registered\n");
}
"""

CPP_DECL = r"""
#include <torch/extension.h>
#include <pybind11/pybind11.h>
#include <pybind11/stl.h>
namespace py = pybind11;
void hip_load_and_register(py::bytes hsaco_bytes,
                           std::vector<std::tuple<int, int, int, std::string, int, int, int, int>> shape_list);
torch::Tensor hip_dispatch(py::tuple data);
void hip_set_fallback(py::function fn);
"""


# =========================================================================
# Build system
# =========================================================================

def _load_module():
    name = "hip_module_dispatch_v3_ext"
    bdir = os.path.join(tempfile.gettempdir(), name)
    os.makedirs(bdir, exist_ok=True)

    arch = os.environ.get("PYTORCH_ROCM_ARCH", "gfx950")
    cuda_flags = ["-O3", "-std=c++17", f"--offload-arch={arch}"]

    return load_inline(
        name=name,
        cpp_sources=CPP_DECL,
        cuda_sources=HIP_DISPATCH_SRC,
        functions=["hip_load_and_register", "hip_dispatch", "hip_set_fallback"],
        with_cuda=True,
        extra_cflags=["-O3"],
        extra_cuda_cflags=cuda_flags,
        extra_ldflags=[],
        build_directory=bdir,
        verbose=bool(int(os.environ.get("INLINE_HIP_VERBOSE", "0"))),
    )


# =========================================================================
# Module-level init
# =========================================================================


print("[hip_mod_v2] Compiling dispatch extension...", file=sys.stderr)
_mod = _load_module()

print("[hip_mod_v2] Loading kernels and registering shapes...", file=sys.stderr)
_mod.hip_load_and_register(_hsaco_bytes, [
        (4, 2880, 512, "fused_m4_m4_n2880_k512", 90, 256, 4, 0),
])


# =========================================================================
# aiter fallback for unsupported shapes (no-op when disabled)
# =========================================================================

_FALLBACK_ENABLED = True

if _FALLBACK_ENABLED:
    def _aiter_fallback(data):
        from aiter import dtypes
        import aiter
        from aiter.ops.triton.quant import dynamic_mxfp4_quant
        from aiter.utility.fp4_utils import e8m0_shuffle

        A, B, B_q, B_shuffle, B_scale_sh = data
        A = A.contiguous()
        A_fp4, A_scale = dynamic_mxfp4_quant(A)
        A_scale_sh = e8m0_shuffle(A_scale)
        A_q = A_fp4.view(dtypes.fp4x2)
        A_scale_sh = A_scale_sh.view(dtypes.fp8_e8m0)
        return aiter.gemm_a4w4(
            A_q, B_shuffle, A_scale_sh, B_scale_sh,
            dtype=dtypes.bf16, bpreshuffle=True,
        )

    _mod.hip_set_fallback(_aiter_fallback)


# =========================================================================
# Entry point — direct C++ binding, zero Python in hot path
# =========================================================================

_mod.hip_load_and_register(_hsaco_bytes_1, [
    (16, 2112, 7168, "m16_mk11b_m16_n2112_k7168", 231, 256, 16, 0),
])
_mod.hip_load_and_register(_hsaco_bytes_2, [
    (256, 3072, 1536, "m256_v6a_kernel", 384, 256, 256, 0),
])
_mod.hip_load_and_register(_hsaco_bytes_3, [
    (32, 4096, 512, "m32_mk4l_m32_n4096_k512", 256, 256, 32, 0),
])
_mod.hip_load_and_register(_hsaco_bytes_4, [
    (32, 2880, 512, "m32_mk4_m32_n2880_k512", 180, 256, 32, 0),
])
_mod.hip_load_and_register(_hsaco_bytes_5, [
    (64, 7168, 2048, "m64_mk3a_m64_n7168_k2048", 224, 256, 64, 0),
])

custom_kernel = _mod.hip_dispatch