<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/61051>61051</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
Serialized PTX constant array with aligned vectors is missing padding
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
stygianguest
</td>
</tr>
</table>
<pre>
The following example fails ([here on godbolt](https://godbolt.org/z/YG3Yxa483)) because the constant array of aligned 3-element vectors doesn't have padding between the array's elements, but at the end.
```
typedef __align__(16) float __attribute((__vector_size__(12))) f3;
__constant__ f3 arr[] = {{1,2,3},{4,5,6},{7,8,9}};
__global__ void foo_kernel()
{
printf("%f, %f, %f\n", arr[1][0], arr[1][1], arr[1][2]);
}
```
This will print `5,6,7` rather than the expected `4,5,6`.
Loading does take into account the padding, so loading `arr[1]` uses an offset of 16 (see godbolt link above), but the array is padded at the end rather than in-between the array elements:
```
.visible .const .align 16 .b8 arr[48] = {0, 0, 128, 63, 0, 0, 0, 64, 0, 0, 64, 64, 0, 0, 128, 64, 0, 0, 160, 64, 0, 0, 192, 64, 0, 0, 224, 64, 0, 0, 0, 65, 0, 0, 16, 65, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
```
The following patch fixes this issue:
[0001-Add-padding-to-vectors-when-serializing-global-const.patch](https://github.com/llvm/llvm-project/files/10850971/0001-Add-padding-to-vectors-when-serializing-global-const.patch)
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJykVU2PozgQ_TXOpRRkiu8Dhw5R9rKHlXYOOydkoAjeduMIm6S7f_3KhnTSmVxGK0WOKNvlV-89u4Qx8jgSlSzZsWS_EbMd9FQa-3GUYjzOZOym0d1H-WMg6LVS-iLHI9C7eDspgl5IZYBhzpLdQBOBHuGou0Yry5I9w3yw9mRY9MLwwPCwTgV6OjI8fDI8_Pwj-vku4jxiWDAsoKFWzIbADgStHo0VowUxTeIDdA9CObAdRFtS9EajhTO1Vk8GOk1mZJhZGMSZ4CS6zuFsyF6IRp_OZ2GYGVg3G4YVNLMFYf0CGruA8T3jLyzl689_2o8TddRDXXsAdc0wD1MHt1daWBe3dpLNbMlRgXldL7hqIz9pWY5LgX5TxKLdepAf6_paal1DHzmkixzAoj2wbMeyXciwQoZVxLI9w4plu5hhlTCs0q9IxrDKGVaFi2T7X045Kt0IVddw1rKDXuv6laaRlAddrGuzdROcJjna3s8hw6R3bH37T6rRT1Ur4NBJnuy4V_4hGD4Log8WN5zZ_in_PwZp4CKVWjABS_lSOFYZSzlMwg40gR3EojS9n6i11LmFN5ZSHtzT8acW3iLOOWDFK4EcrQbRtnoeF0OsLnK4jQa1bmApv6si5TAbMiBG0H1vyDqfhqm7E4boehlAyfEVRKPP5E2w-O7LlSCNP4y6Oy9-K0uO21-8fPNx9PKUt-AsjWwUQeD9BYG3r0MXNPkqRZzf-Yw7ZH4InewVpNFX5Dak8bPvx-g1w2M4fZojLPBZGPFp7iVF8pj5WfQ3h_tb853MqxXvn8GTsO0AvXx3FnImlcbMdJMj2XHOw-1L121XL22t3q5v1vYy0Lg1NEmh5KebW-7n1osV-NxPH1Fph7kJWv3G8KDU-fq3PU36X2otw0MvFRmGh5DnCS-ykOHh_wK5vg-broy6IirEhsowzbKkwCSPNkNJqeDYtdg2vEXMiz5JKIxF3PZtR3EcbmSJHCOOmIdRkmIUCB52aRKnXZjmWR63LOb0JqQKXDWuR2w8nWUa8iTcKNGQMr5PIY50WblG94hsptIz0MxHw2KupLHmlsVKq6j8e62POvjrxz-PzeUi7fDVXq5NRRp4k8YsSnvWNvOkyt-WwyN1evhK_gsAAP__wOYrLg">