<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=http://email.email.llvm.org/c/eJytWMty47YS_Rp5gzJLIm09Fl7YclyZxUymKk4mWblAEpIwJgkFACV7vj6nGyBFSbZzUzdTGloiHv06fbqB3JSvN6P0Fp9f27qW9jX8GI3vR-P4_CQ2cqeEFMvf7m_Fs7KNqoTfSC-8fFYOA6PpuG20afBXSCdMo4RZCe2d2Eora-WVdYn4Jbz2GzVcMEpnTtSqzpUV_nWL_XTc0nnbFp72HKXzQlaVKun1V28dvdSNqF9FrRtdy0rsjX3WzVqoF1lvKzVKF0HFwjRe6sYJZ2oFdcqSZpWtEt6wKhBt7KuQlV43tWq8sOqvVltF311nhtk3UUfY8QmitddY8UOdWkPm8wK8t0OzyAaM30svo_K08vO3n0hT2ZRQzdHSTq3o5kq2TbGBzMbpUg1HNmavdsqO0mUwLXqW7O00H8iT262SljbPlfihrIEvTevF5aVojJd59YrVxtF6a5Xbmob9FJXp3GbVGmbGrWMcoJsXTqnaiUo_K7GsJGZWxjwjjP5ouQ5GmP8RB-QVOKWt8Z1jCQh0IaNIBFhuTEWYCf50koJcyUK5ZAjhbxvVwLR6qys2C1MHYCYfBrXdxrRVKWrgGn4J0reebJY5uWu_0cVmaMApcIsW7ms83Om8sXAyFDvOD7YEcV4rBEmSlULadUtwIz3ylkeLCuHCLqWJjtCeHUQy99ZA1YFubDwJl04l4pvq8In1K9NiMY_mLUWAAPIdmOKFe-03YnR95ze2dX50fQ-QbrzfulEGEnjAZ40ZbZ7AdfjR7HSpJb7E-RG52KBGnj_Toux2J62WzT9vFpakD938dDE0H14BUE3ulN3Bi6wo4gr9OW6FoWQwhtOtT6WjkAce-x04gaXuLVazamtN2RZq4CAWVLeV1yCRCIpd3AMa-b0CkGbJOBmz6ZMr-nppiwkpD2X2ChiVIU4RC0CTAwGK3EpO5S_Ri4hmsyuKCEtM4FgjG8FhG53DC64F2OAJjVCrI-N-9dJ6QjLr-4GXq2rX_bmEtd8VCDV9wFhNgHoox9PJtVxkeT4tZ1KVxVWWpVlZzBfXs5W6nqn5VJVzqcg6kC183NEjUpO414DkTKlXupCMKBAGJ49rrYo1Ap7oAA5bt69iL10D870wgG8NEmUqSsQjBSG-C7vp4I-1apRlOG-NczqvmE1Vsk6QkIhH60JSR1AGSkQoCEa03iospe1AIUUBCkJJ4OTQLx0bWIUacu7nAJvP79WYU0wtQyhjXk3HBeFnlN7hIy5_ycSl8-Uouy_Cq8kV6LdoS3m53raX0hYbjLn6aTYW9V4lRSsuTf-toonWM-E-QmHWlfNEVs6cJQuRFeZCKK-4Lbl0Xt4_fPrj6dMj85AhCkIUQ5ntV0WyJiciswLhEbGwm4LUHqimAMaPHQYi5w-p2_kw001RtcjYUbaU1dpYiKpH2U9vjWsDfZR8b3gHABt7GHxjCklONidTopFcs8Lr2V34IkSxQWV0nMEu5WfGzyt-XvNzys_ZKLs7iFyVaiWiQ4-2mg-nKVTSVTe-M7okBnaITnorSqrO_dzZ_eE7P70CzKQPVlXUHzySGi3K6LpBoL_0RkbzuNqfmvcYxFzffSFO_i-kcS1jYY-oe2cCWYtsSatpjSh7QYIDILYfaPH0tK5MLqunp-AsD9Z0K2PBYvNeYrbstCINr0hIW8vtkbrN0S_yNvoaEDsRxBkCOCL9ZA0X3BOfKFl-Kl-SlxMdhUBE03kjSPD4ZYV_VAsHGwrBXHwbOK2gbqonNOraYFarunpHidWoF4_-CY2Faog7Y8dDjQMZRh0j7bPV3FIgd_dIIklMGPrX5CCZbQ2Va5KJj4vwf1YeFifyoWNXW0FRR6TOhu016iQ3yXsNbkGJfB4QzJ5KBhjHvsYmtISRCJEFPE8sJaP6F-SrpEwi3jXwzpEcDwBIiDsPJUV8SZBZHOYNghmRM9jyXNKbMo6F5caggVVUpOZvwZCrA_dQkeUOMAfQCAezuzT7jXATUBeQhwxI-eUEY7OhIv-YMNlp0p5NGyZvrx-VcdhAPshxJGiiPdQk4BUY7_BilE57N425KVwMdgxuZI-WZ0l2kpXoM3giLXHoGqKM01UfO3HztJNVq9APzifTt5YfMUaYy1Kbtqq23g50pzrzGX2EKViTabfxshcyUJO3o5-G4Hbw8uLYHwe2y5bhM3lnQyLj8MF2kfvemobXnWKLU-VVXXAc-3Ucp7N1_9ag5WD3e7XThXo0P5vQmh1r8GAVbXeuYDeJM2ZtYuSZNA87wE-UumgrdZ_mImQysEI_Qp-l30vrAZ6pXWJ_HxByltbEFWLgrp4P0knEZbdmSIdHEt-ROUrTcKrDlyOB4mD8Cn2eOhrs2eXdffuXSMnqIxK0yre2YWFHtXkwSfNJps_1U-46i9WB5rIjJmVBGIkzH8QYz1sxOQp_L7prJ8-b8lvncPau3rivejR0H4HWlQ7adEuw3aomtrTxsNWdGEOnTie5YdM86JJpiNZwOcKaL79_ffyDzt5BNhQ_qa6mzE3lE2Mh7uEH_n-bff1z8VfmvotE_GlabtlJvXAPEM6ulPFoCqjN5r67O_J4Ip_puCoTvkfrFOK7sQTEI6twlxTaPz6jxksv7jE21E14HNY4Ud-8c1lSh1-pThH0fz4qQlq4Z7118WIof_X9PYzt73SSi_ImKxfZQl7IFhPtTf3cAskXra1u_nXjwdUfRPBwnc0m44vNzWp2JaeT-WQ-zrJsMVFlin4km5bX01RO8zy9qGSuKneDxEPqNGofGghKo-v7i_9fA32TjlN8JuPJ1Xg8yZJVcTUvFmoxnUyzdHGlRldjhayoEtqHon5hb3jLvF07DFbaeXcYBHCIJxUrDA299pW6oSsNghKHhS804h1fd24Op63Y_hEcK-XpTAs-ihdQ_TVnjMsFG3LDVvwN6ZG_2g>53710</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            CUDA: unions in kernel arguments not copied completely if member contains padding
        </td>
    </tr>

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

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

    <tr>
      <th>Reporter</th>
      <td>
          mkuron
      </td>
    </tr>
</table>

<pre>
    **Summary**

I have a CUDA kernel that takes a `union` as one of its parameters. One of the `union`'s member types is a `struct` (called `Ptrs` in my minimal working example) that contains some padding due to the memory alignment requirements of its own members. I initialize the `union` as its other member type (`Data` in the MWE) and pass it to the kernel launch. Inside the kernel however, some of the contents of `Data` appear to be zeroed out -- notably those corresponding to the padding region of `Ptrs`. It seems like Clang looks at the padding inside one of the `union`'s member types and assumes that all the members have holes in the same places.

When compiling the CUDA kernel, Clang should make no assumption about which one of the member types is currently stored in a `union` that it gets as an argument, but it clearly does and it's the wrong assumption in this case. We initially found this bug in conjunction with [thrust](https://github.com/nvidia/thrust) and [mpark::variant](https://github.com/mpark/variant), but it can be observed with plain CUDA code too as in the MWE.

**Versions**

I reproduced this bug with multiple Clang versions between 7.0.0 and 14.0.0-rc1, as well as the current master branch. Nvidia's nvcc compiler does not exhibit such an issue.

Starting with https://github.com/llvm/llvm-project/commit/d0615a93bb6d7aedc43323dc8957fe57e86ed8ae, my MWE required a minor modification to make sure that the argument copy wasn't optimized out. This optimization is not generally possible (e.g. when using thrust), so we cannot rely on it accidentally fixing the real issue.

**Minimal working example**

Compile with `clang++ -O3 -std=c++14 --cuda-gpu-arch=sm_70 mwe.cu -o mwe.cu -lcudart`. The issue can also be observed without `-O3`. Add `-DFIX_IT` to get a struct without padding, in which case the issue does not occur.

```cuda
#include <algorithm>
#include <iostream>
#include <vector>

#include <cuda.h>

struct Ptrs
{
  char s1, s2, s3, s4, s5, s6, s7;
#ifdef FIX_IT
  char s8;
#endif
  void const * data;
};

template <class T, unsigned N>
struct Data
{
  T data[N];
};

template <class T, unsigned N>
union DataType
{
  Data<T,N> d;
  Ptrs p;
};

__global__ void transform(DataType<unsigned, 4> umap, unsigned n, unsigned * result)
{
  const unsigned i = threadIdx.x;

  if(n > 0xffff) {
    /* This condition is never true, but the next line ensures that `umap` is copied into writable memory.
     * Clang 13 (https://github.com/llvm/llvm-project/commit/d0615a93bb6d7aedc43323dc8957fe57e86ed8ae)
     * introduced an optimization that will otherwise mask the issue we are trying to demonstrate.
     */
    umap.d.data[i] = 0;
  }

  if(i < n)
  {
    result[i] = umap.d.data[i];
  }
}

bool test()
{
  std::vector<unsigned> map{23U, 0xffffffff, 42U, 13U};
  DataType<unsigned, 4> umap;
  Data<unsigned, 4> d;
  std::copy(map.begin(), map.end(), &d.data[0]);
  umap.d = d;

  const unsigned mod = map.size();

  std::vector<unsigned> h_values(16);

  unsigned * values = nullptr;
  cudaMalloc(&values, h_values.size() * sizeof(unsigned));
  transform<<<1, h_values.size()>>>(umap, h_values.size(), values);
  cudaMemcpy(h_values.data(), values, h_values.size() * sizeof(unsigned), cudaMemcpyDeviceToHost);
  cudaFree(values);

  bool good = true;
  for(int i = 0; i < mod; ++i)
  {
    std::cout << h_values[i];
    if (h_values[i] != map[i])
    {
      std::cout << " wrong";
      good = false;
    }
    std::cout << std::endl;
  }

  return good;
}

int main()
{
    bool good = test();
    return (good ? 0 : 1);
}
```

**Assembly**

To see what is happening, compile the MWE with and without `-DFIX_IT` and compare the NVPTX assembly: https://godbolt.org/z/W7PY9q3sj . You can see that in the former case, it generates `ld.param` and `st.local` instructions that cover the entire size of the `union`, while in the latter case it skips those bytes that are padding.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJytWMty2zYU_Rp5gzGHIvVceGHL9TSLpJmp27QrD0hCEmKQUAFQsvP1PRcAKUqy3XamGYamiMd9nXvuBQtdvd6Msltcv7Z1zc1r-DFK70dpvH9iW74XjLPVb_e37FmYRijmttwxx5-FxcBolraN1A3-Mm6ZbgTTayadZTtueC2cMDZhv4TXbiuGC0bZ3LJa1IUwzL3usJ-MW1pn2tLRnqNsUXKlREWvvzpj6aVsWP3KatnImit20OZZNhsmXni9U2KULYOKpW4cl41lVtcC6lQVzapawZz2qkC0Nq-MK7lpatE4ZsRfrTSCnm1nhj40UUfY8QmipZNY8UOcW0Pm-wV4b4ZmkQ0Yv-eOR-Vp5edvP5GmvKmgmqWlnVrRzYq3TbmFzMbKSgxHtvog9sKMslUwLXqW7O00H8jju53ghjYvBPshjIYvdevY9TVrtOOFesVqbWm9McLudOP9FJXp3GbEBmbGrWMcoJtjVojaMiWfBVspjplK62eE0Z0sl8EI_S9xQF6BU9oazz6WgEAXMopEgOVWK8JM8KflFGTFS2GTIYS_bUUD0-qdVN4sTB2AmXwY1LZb3aqK1cA1_BKk7xzZzAty12Ery-3QgHPgli3c1zi40zpt4GQodpof3hLEeSMQJE5WMm42LcGN9ChaP1oqhAu7VDo6QjrvIJJ5MBqqDnTzxpNwbkXCvokOn1i_1i0W-9GipQgQQL4DU37hQbotG03v3Na01o2m9wDp1rmdHeUggQdcG8xoiwSuw49mLyvJ8RDnR-Rigxp5_kyL8ts9N5I3_7xZWJI9dPOz5dB8eAVA1YUVZg8vekURV-jv41ZqSgatfbr1qXQS8sBjvwMnsNS-xWpG7Iyu2lIMHOQF1a1yEiQSQbGPe0AjdxAA0jxJk9SbPp7Q47Upx6Q8lDkIYJSHOEUsAE0WBMgKw30qf4leRDSbfVlGWGKCjzWyERy2lQW8YFuADZ6QCLU4Me5Xx40jJHt9P_CyUvvuzzWs_S5AqNkDxmoC1EOVzsZTvsyLYlbNuajKSZ5neVUultP5WkznYjET1YILsg5kCx939IjUJO7VIDldybUsuUcUCMMnj22NiDUCnugADlt3r-zAbQPzHdOAbw0S9VSUsEcKQnwXdpPBHxvRCOPhvNPWykJ5NhXJJkFCIh6tDUkdQRkoEaEgGNF6I7CUtgOFlCUoCCXBJ4d86djACNSQSz8H2Hx-r8acY2oVQhnzapaWhJ9RdoeLXf-Ss2vrqlF-X4ZX4wnot2wrfr3ZtdfclFuM2fppnrL6IJKyZde6f1I00ThPuI9Q2Ovq84Qrqy-ShcgKcyHUr7itfOm8vn_49MfTp0fPQ5ooCFEMZbZfFcmanIjMCoRHxOLdFKT2QNUlMH7qMBC5v0jdzoe5bErVImNH-YqrjTYQVY_yn94alxr6CP7e8B4A1uY4-MYUkpxsz6ZEI33NCq_nd-GBsXKLymh9BtvM33N_n_j71N9n_j4f5XdHketKrFl06MlWi-E0gUq67sb3WlbEwBbRyW5ZRdW5nzu_Pz77uxOAGXfBKkX9wSOp0aKMbhoE-ktvZDTPV_tz8x6DmOndF-Lk_0Oar2Ve2CPq3oVAr0W-otW0hlW9IOYDwHYfaPH0tFG64OrpKTjLgTXtWhuw2KKXmK86rUjDCQlpa747Ubc5-UXeRl8DYieCuECAj0g_WcIF98Qnglefqpfk5UxHxhDRbNEwEpy-rPGPauFgQ8Y8F98GTiupm-oJjbo2mNWKrt5RYjXixaF_QmMhGuLO2PFQ40CGUcdI--ykbymQuwckEScmDP1rcpTsbQ2Va5yzj4vw_1YelmfyoWNXW0FRJ6TuDTtI1EnfJB8kuAUl8nlAMAcqGWAc8xqb0ApGIkQG8DyzlIzqX5CvkiqJeJfAu49kOgAgIe4ylBTxFUFmeZw3CGZEzmDLS0lvyjgVVmiNBlZQkVq8BUNfHXwPFVnuCHMAjXAwv8vy3wg3AXUBeciAzL8cY2w-VOQfEyY_T9qLacPk7fWjMg4byAcFjgRNtIeaBLwC4x1fjLJZ76bUN4XLwY7Bjd6j1UWSnWUl-gw_kZZYdA1Rxvmqj524fdpz1Qr0g4vx7K3lJ4wR5nqpTavUzpmB7lRnPqOP0KXXZNZtvOqFDNT029FPTXA7enl56o8j2-WrcI3f2ZDIOFzYLnLfW9PwulNsea68qEsfx36dj9PFuv9q0Gqw-73Yy1I86p91aM1ONXgwgra7VLCb5DNmo2PkPWked4CfKHXRVso-zVnIZGCFfoQ-S76X1gM8U7vk_X1EyEVaE1ewgbt6PsjGEZfdmiEdnkh8R-Yoy8KpDg8nAtnR-DX6PHEy2LPLu_v2L5GS6iMSNMK1pvHCTmrzYJL0J5k-18-56yJWR5rLT5jUC8JInPnAUtxv2fgk_L3orp28bMpvrcXZW73xvepR0_cItK500KavBLudaGJLGw9b3YkxdOp0khs2zYMumYZojS9HWPPl96-Pf9DZO8iG4mfVVVeFVi7RBuIefuD_t_nXP5d_5fY7S9ifuvUtO6kXvgOEsytlPJoCarN9390deRyRzyxVVeK_o3UK-W9jCYiHq_AtKbR__owaP3r5HmNL3YTDYc0n6pvfXFbU4SvRKYL-z0VFSAv7LHc2fhgqXl3_Hcb033SSq-omr5b5kl856ZS4oeM5ucWL8Ifz-L2qOwOGk0NsZci1Sjg6nyG34seU_pNdlHHVGnXzn5sY30mAVB6m-XycXm1vFkuRjudFPinKIhOLMi3T-TpPp9l6XM2zSlwpXghlb5DESMNGHEIzQik5vb-SN1ma4Rqn40majvNkXU4W5VIsZ-NZni0nYjRJBfJDJaQHxf_K3HiVinZjMaikdfY4CAgRYwovDvvzFm42N_VzCx648pJvvOZ_AyzXqyw">