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

    <tr>
        <th>Summary</th>
        <td>
            Cuda memory error when using RDC and a constexpr function
        </td>
    </tr>

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

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

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

<pre>
    The code below fails with a memory error if I compile and run with RDC.

```
#include <iostream>
#include <cassert>
#include <cstdint>

#include <cuda_runtime.h>

static constexpr int bucket_group_size = 255;

class element {
public:
  static constexpr size_t num_limbs_v = 5;

  element() noexcept = default;

  constexpr element(uint64_t x1, uint64_t x2, uint64_t x3, uint64_t x4, uint64_t x5) noexcept
      : data_{x1, x2, x3, x4, x5} {}

  const uint64_t* data() const noexcept { return data_; }
private:
  uint64_t data_[num_limbs_v];
};

struct element_p3 {
  element_p3() noexcept = default;

  constexpr element_p3(const element& X, const element& Y, const element& Z,
                       const element& T) noexcept
      : X{X}, Y{Y}, Z{Z}, T{T} {}

  element X;
 element Y;
  element Z;
  element T;

  static constexpr element_p3 identity() noexcept {
    element zero = {0, 0, 0, 0, 0};
 element one = {1, 0, 0, 0, 0};
    return element_p3{
 zero,
        one,
        one,
        zero,
    };
 }
};

template <class T> __global__ void bucket_accumulate(T* bucket_sums) {
  for (int i = 0; i < bucket_group_size; ++i) {
    if constexpr (true) { // change to false and things work
      bucket_sums[i] = T::identity();
    } else {
      auto identity = T::identity();
 bucket_sums[i] = identity;
    }
  }
}

int main() {
 using T = element_p3;

  T* data;
  auto rcode = cudaMalloc(&data, sizeof(T) * bucket_group_size);
  if (rcode != cudaSuccess) {
 std::cerr << "cudaFree failed: " << cudaGetErrorString(rcode) << std::endl;
    return -1;
  }
  bucket_accumulate<T><<<1, 1>>>(data);
  T res[bucket_group_size];
  rcode = cudaMemcpy(res, data, sizeof(T) * bucket_group_size, cudaMemcpyDeviceToHost);
  if (rcode != cudaSuccess) {
    std::cerr << "cudaFree failed: " << cudaGetErrorString(rcode) << std::endl;
    return -1;
  }
 std::cout << *res[0].X.data() << "\n";
  rcode = cudaFree(data);
 if (rcode != cudaSuccess) {
    std::cerr << "cudaFree failed: " << cudaGetErrorString(rcode) << std::endl;
    return -1;
  }
  return 0;
}
```

The error I get is 
```
cudaFree failed: an illegal memory access was encountered
```

If I change
```
 if constexpr (true) { // change to false and things work
 bucket_sums[i] = T::identity();
    } else {
      auto identity = T::identity();
      bucket_sums[i] = identity;
    }
 }
```
to
```
    if constexpr (false) { 
      bucket_sums[i] = T::identity();
    } else {
      auto identity = T::identity();
 bucket_sums[i] = identity;
    }
  }
```
the code works.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcV0uT2joW_jVicyqULWNML1i0TXomi9kkLLp7Qwn7AJrIEqVHdye_fkryE0wnmXtv1X1QFNjSeen7js6RmDH8KBHXJM1JupkxZ09Kr7XcOy1ne1V9W29PCKWqEPYo1CscGBcGXrk9AYMaa6W_AWqtNPADfIJS1WcuEJisQDvZCH7eFHMSbUh03_4uo_bbvNKEy1K4CoEkBVfGamQ1ST7emi6ZMajte7PGVlyOZm_JuIrttJOW1zg_XYkayywvoVTSWHw7a-DSwt6VX9Hujlq5887w797MBmiakiQfK5eCGQMosEZpgWTt5NntBS9J0koBTJx4mzsL0tU7weu92b0ED9f2obNN6IrQO5AK30o82yBc4YE5YScqg5dB2XFpl4udhbeY0AKGV3r5mly-Li5f03EMnTv_Ick9VMyyHcnyxkNjuLHXmHlLSbYJGGWbGwH3bghtbLVLbiaHhWc5aLROy9ZhkkNv8Kz5C7M4Ar6PvRFO8xHiJN0M2GWbKxyN1a60HYS7czLQC6PR30VMo9-ssOdqCY8ersno083RZ0KLMROTz0Rj-yMSH0mWP3owaAFPJMuf2udnkuXP7fOWZPn2PS67vfDYr74fehqG-rHnG2PbCXCT7TNihVcoLbffJkQMbA2mv6NWgSKS5ZFfy-RnlAe9lpLYKcU_VQLo8nNEch-LD2DCmJL4S2MT5QvHAxWTXLZYnwWzTTEMJWtLko-w2x2F2jOx28GL4lVX9VhZutp5eUJXW78d2wnjauMxHkF7UBoIXfmayQNGkd-Q_rGYFtGwV2lOaM6vzIDvJQO9hK6sdtgKAaEPhD5AeWLyiGAVHJgwTcexJy6PBl6V_jrGahxwmnOSbkJwW18ZkvvLlLmgzqc1euMXwQEwZ1Wfar9i650IevErp93LBYkjBj3ANeOyzfIhOme4PMI2GB8l3PUO2vZVdXAc1qRDp_favkv-hwmhyuBk2dTgIrQqdWhS4Q5G6TAi9gJFfvAEtoZp3Nn-4soSzVUCGVs1OJaotc8anziEUq_woBHD6QO9jB_tBPzsv9B-9KeQL1ZzeewcBuuNUG8aZSVu7c8P8Wh0xMF0GySF3y-NXZIUoQjEYSR86arBagzCFjR65qdYpeNicQU_1uXZp5LXpQX8XxQUIxMbfOElbtW_lbG_nRyAvyI_Q0jK2SGk-wbviKSb-eN8dIAYgiZpIf3fe_D7Bd0i8--PWTcdXRx4bh7Nw6-_AjRn_E9wRAvcwE3pG0tiErgQeGSiuyywABO8MgMoS-WkRY3VD7x_CveKUO1vSv1xzeLP7xM_alc_aRbvUWjVbdRudNmATY_cP6uBXoHSXWo9-2Y-q9ZJdZfcsRmu4-VdRpN0sVjNTusyjdIsOkSrmC6XaRxnmEZxFMcpjWJGV8mMr2lEkziKE0rjjGbzhLGY7tPDIt1XiwUysoiwZlzMhXip50ofZ9wYh-vlXZQkM8H2KEy4fVMq8RXCZChOm5lee50Pe3c0ZBEJbqwZrFhuBa4LV13dwl9PKNtjwOdNERKdjVg-OFlaruTMabE-WXs2HvewSY7cntx-Xqqa0Afvp_37cNbqv1haQh9CdIbQhxD9_wIAAP__MuacBA">