<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/117562>117562</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU] [MI300] How does kernel initialized with sgpr mapping and how its offset computed ?
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          yiakwy-xpu-ml-framework-team
      </td>
    </tr>
</table>

<pre>
    I checked the assembly ABI and code object specification, and surprisely found that our float, double pointer takes size of 8. While int still take size of 4. I am wondering why we are using a 8 byte to represent a memory address of (hsa kernel dispatch packet address)?

Also I want fully understand offset computing logics for MI300X :

I checked source code, the following instruction means kernel argument pointer address stored in s[0:1], and s[0:1]+0x1cd should point to blockDim.x

.loc    21 275 58 prologue_end          ; /opt/rocm-6.2.0/lib/llvm/bin/../../../include/hip/amd_detail/amd_hip_runtime.h:275:58
s_load_dword s6, s[0:1], 0x1c
s_waitcnt lgkmcnt(0)
s_and_b32 s1, s6, 0xffff
It is obvious the blockDim.x and blockDim.y are fetched in a 32-bit register (dword), and we extract it with 0xffff mask, but how should we determine the offset 0x1c?

I checked HSA specification, but it is all for RDNA. Do have document about how this kernel initialized ?

I also want to know where I can ensure that s[0:1] point to a kernel dispatch patchet and kernel arguments, while for GCN arhictecture, usually s[4:5] point to kernel dispatch packet and s[6:7] point to kernel argument region.


</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJx0Vc2O2zYQfhr6MrAgUdLae_BBm4WTPSQoWhTtbUGRI4s1RQrkcLXO0xek_7Jp6oNgScOZ-X5mJELQB4u4Y-0Ta59XItLo_O6kxXE5rd_nuJ7MevBiwsX545pQTKveqdPuBeSI8ogKaEQQIeDUmxN0Ty8grALpFILr_0FJEGaUetBSkHaW8U85IEQ_ex3QnGBw0aY0gsBFD4NxglKYcrE3CLPTltADiSMGCPo7ghtgW8BfozYI2hIE0sbkgNv7poAXEBMszir02h5gGU-wIAiPEEN6IGAL_YkQyIHH2WNASyBgwsn5EwilPIaQkjG-HYOAI3qLBpQOsyA5wizkEekayPgjq_es7FjZdSY4eIFFWIIhGnOCmNoIlKC7YQhIIN00R0qNGHfQMsDgPHx9qcvyb2B1d050pzm46CVmYhM5ifXBGeOWlEHbQD7KRDBMKGy49ir8IU4J1pXFK6pAzqMCbSGw9qlkdVex9vmmzodnT-V7JRWE0UWjzpkSZ71x8visp-L93GthnGTlI6-Ab1potzB7Z9wh4itaBbcfq5-A8b2bifG9d3JaPxS8KBnfG92nq3mbGN_32jK-L4oPF22liYmA_ahnxvdiUq8KSWhzuRn1_OqjJT1hMbK645uW1V27ZWUXXo0T6lUtzisIDwnqz9ATzhy5CE3SEpjDcZKWGN-WSd70Slj12tccQpUzPJzPDcMwJLkIdADXv2kXQ9boTlIm9nZ7ykYckOR4lkFAzde9JvB40CFJxfg2N5sqX3RZEPCdvJAEmmDRNF5qwyTCMUX1kWB0y1WsBUEhoZ-0xdzPxXwZ6dWtd5N9-aP777imlDoDE8Zkl_7-_K0r4NnBKN4QlJNnj4neXarTqG8W1FaTFkZ_RwU_lBRpRPKAkIOjdQssI3qEF5DCAtoQPZ6Xwo8q3d33q3FMZFIm6if7h4Rjyfsi9f_50zcQftSSUFL0eaBiiCINaqrWJM98qPZ_o3-ZlQdWd5tfHbjNX1LV2eKMfqV2tXqsH8UKd9Wm5m1TN5tqNe6UanvEatNu2q1qSl73G9k8bodW9Q-qkWqld7zkTVXxtmrKpq6Kx03PRd02PS-3pWgla0qchDZFGqPC-cNKhxBxV1Wb9oGvjOjRhLztObe4QH7LOE_L3-_SoXUfD4E1pdGBwj0NaTL5M9F9ff78258JLGuf8r5K_7-4BZTDX6qefRoOs4dJzHNevVZln2gKH9fh2SOr6M1uJJpDWoR8z_j-oGmMfSHddF8RudvZu_SVScshQQmM7y9Y33b83wAAAP__ciVGRw">