<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/59156>59156</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
Incorrect code generated when -O1 or more is enabled
</td>
</tr>
<tr>
<th>Labels</th>
<td>
bug,
clang,
cuda
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
carlosgalvezp
</td>
</tr>
</table>
<pre>
Hi,
We notice the following minimal example code produces incorrect result when using -O1 or later:
```cpp
#include <cstdio>
struct Member
{
int value;
bool valid{}; // <<<< Removing {} makes it work
};
class Wrapper
{
public:
__host__ __device__ Wrapper(){}
__host__ __device__ explicit Wrapper(Member member) :
member_(member)
{}
Member member_{};
};
__device__ inline Member makeMember(int const input) //<<<<< Removing inline makes it work
{
return Member{input, true};
}
__global__ void testKernel(int const input, Wrapper* const d_wrapper)
{
auto const x = makeMember(input);
*d_wrapper = Wrapper(x);
printf("Wrapper member value: %d\n", d_wrapper->member_.value);
printf("Member value: %d\n", x.value);
}
int main()
{
Wrapper* d_wrapper{};
cudaMalloc(&d_wrapper, sizeof(Wrapper));
testKernel<<<1, 1>>>(123, d_wrapper);
cudaDeviceSynchronize();
}
```
Compiled with:
```
clang --cuda-path=/usr/local/cuda-11.7 -O3 -std=c++14 --cuda-gpu-arch=sm_75 main.cu -L /usr/local/cuda-11.7/lib64/ -lcudart_static -ldl -lpthread -lrt
```
Expected result:
```
Wrapper member value: 123
Member value: 123
```
Actual result:
```
Wrapper member value: 0
Member value: 123
```
NVCC produces correct results, on the other hand. Is this a compiler bug, or some hidden bug in our code?
Thanks!
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJylVduO4jgQ_ZrwYoFyIZA85IGGae1ot3el3dH2I3JsQzzt2JEvNDNfv2UnhEBfHnZQFBK7fOrUqUpVreiP6jcepdso3kXxpr8_MySV5YQh2zB0UEKoVy6PqOWSt1ggdsZtJxgiijLUaUUdYQZxSZTWjFikmXHCoteGSeSMPzn_K0FKI4Et01G2mTqLVnF_ka4bVtIMsIQD8CjbEmMpV1H2ZXrIWO3A0RNrawDsN9YP_QOCH5cWnbBwLMomq7VSwi9z6o3XO9hEUfoIl3c0Xuhv1qqTp92boRa_-PggIqVfLt52I3R_JwIbg5417rq3lDpXC07GyD2Z_b5Rxu738EDZCcSGx8vptIjScuD46QF27gAXmF1P9pqgtpcmLdGNV__rt_ZgOxpdDe693sDtR90-UmHCjUvBJRsBQMMhXWnh00OUNBaMOmc9yyER0zzcpGJAezcVkxRrZp2Wl8JYPwz4WwQFw-6p3_A-ClVjAbxPilNkmbG_My2ZeI_u9qr3Ztih-9fLUvkOL-ysGizPkJHdvR69CjfVCtgjaDhyzfH53rbTwPEQ6iYdzIaUXT6DDeDlNMq3Ekx8BCP2HL6tIb2L3vgT8KfPQc9vEe6E9lq2mMuhxt8qNVH2qult2Xkz4ih-wtCZSEBaTfTfIsN_MuUZP19zchfUJMGXYkv80cR3mv5KiyTNbqS6B_EkdqHe__khSaOVBMdDZG_jvzS6qRxb1XZcMIpeuW0-6oxjg_GddO6dzjvszXfwxTgDtB5BBwyl-hg2k2SxhpaboTn0TrACiR7gSpaX08fOzbEmHsG0-3UeMrIgDs3_QB9D-jVer5a-X86FX9Z2byyGSQHvVMCts41mmMKTtp9E_eXcwZyAqPtJ8XncH9WzT04wuK_JceNd3xtiHcywX_Ec_x-_f_673V6H5e2oNL7KlAzjVsFNowZLukBfDSxxgzDYhzrRqHbHYKyRUS1DDacUpiysQm9CyukwlKMMJtrE9zeAewEnyYxVyWpVpklapPGMVhktsxLPLLeCVV_HAR4G-5FJprHPUhjkwwhvlWYIKDGJa6jbmdOiaqztjJcytPAjVLKrF8DYF4w4Xf7mEPx3QIdXboxjwOcxL5N8NWuqLE8zzDKar3FeJMuSFfhwOJRFTJMsPxTrmcA1E6aKcijkNGgQ-g38hc9i8gpl6d_y3YxXaZymCaQlWSZZVi7iJVnVBcHFMiEHEmfRMmZQ-WLh6S2UPs50FZiCAwObghtrrpsw3_lRMhZYAD709EbpimAtlDlicWI_u1kIrQpx_QfTYb5E">