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

    <tr>
        <th>Summary</th>
        <td>
            Clang cuda functions not handling concepts correctly.
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            clang
      </td>
    </tr>

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

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

<pre>
    Take the example code posted below.

If I compile, I get this error
```
> clang++ -x cuda -c --cuda-gpu-arch=sm_70 -std=c++20 main.cc
clang++: warning: CUDA version is newer than the latest partially supported version 12.1 [-Wunknown-cuda-version]
root@3aa37681f2bd:/src/concept_bug# clang++ -x cuda -c --cuda-gpu-arch=sm_70 -std=c++20 main.cc
clang++: warning: CUDA version is newer than the latest partially supported version 12.1 [-Wunknown-cuda-version]
main.cc:116:3: error: no matching function for call to 'kernel'
  116 |   kernel<<<1, 1>>>(bucket_sums, generators, scalars, 0);
      |   ^~~~~~
main.cc:26:17: note: candidate template ignored: constraints not satisfied [with T = E]
   26 | __global__ void kernel(T* bucket_sums, const T* generators, const uint8_t* scalars,
      |                 ^
main.cc:25:11: note: because 'sxt::bascrv::element97' does not satisfy 'element'
   25 | template <bascrv::element T>
      |           ^
main.cc:13:3: note: because 'double_element(res, e)' would be invalid: no matching function for call to 'double_element'
   13 |   double_element(res, e);
      |   ^
1 error generated when compiling for sm_70.
```

But the concept element should be satified. In fact, if I uncomment the line
```
  // f(bucket_sums, generators, scalars, 0); // uncomment this line and things compile. ¯\_(ツ)_/
```
the compilation will succeed.

Here is the version of clang I'm using
```
clang++ --version
Ubuntu clang version 18.0.0 (++20230913042131+f3fdc967a87d-1~exp1~20230913042254.1182)
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin
```


code:
```
#include <iostream>

// element.h
#include <cstdint>
#include <concepts>

namespace sxt::bascrv {
//--------------------------------------------------------------------------------------------------
// element
//--------------------------------------------------------------------------------------------------
template <class T>
concept element = requires(T& res, const T& e) {
  double_element(res, e);
  add(res, e, e);
  neg(res, e);
  add_inplace(res, res);
  { T::identity() } noexcept -> std::same_as<T>;
  mark(res);
  { is_marked(e) } noexcept -> std::same_as<bool>;
};
} // namespace sxt::bascrv

// kernel.h
namespace sxt::mtxbk {
template <bascrv::element T>
__global__ void kernel(T* bucket_sums, const T* generators, const uint8_t* scalars,
                       unsigned length) {
  (void)bucket_sums;
  (void)generators;
  (void)scalars;
  (void)length;
}

template <bascrv::element T>
void f(T* bucket_sums, const T* generators, const uint8_t* scalars, unsigned length) {
  (void)bucket_sums;
  (void)generators;
  (void)scalars;
  (void)length;
}
} // namespace sxt::mtxbk

// example_element.h
#include <cstdint>

namespace sxt::bascrv {
//--------------------------------------------------------------------------------------------------
// element97
//--------------------------------------------------------------------------------------------------
struct element97 {
  uint32_t value;
  bool marked = false;

  static element97 identity() noexcept {
    return {0};
  }

  bool operator==(const element97&) const noexcept = default;
};

//--------------------------------------------------------------------------------------------------
// double_element
//--------------------------------------------------------------------------------------------------
inline void double_element(element97& res, const element97& e) noexcept {
  res.value = (e.value + e.value) % 97u;
}

//--------------------------------------------------------------------------------------------------
// neg
//--------------------------------------------------------------------------------------------------
inline void neg(element97& res, const element97& e) noexcept {
  res.value = (97u - e.value) % 97u;
}

//--------------------------------------------------------------------------------------------------
// add
//--------------------------------------------------------------------------------------------------
inline void add(element97& res, const element97& x, const element97& y) noexcept {
  res.value = (x.value + y.value) % 97u;
}

//--------------------------------------------------------------------------------------------------
// add_inplace
//--------------------------------------------------------------------------------------------------
inline void add_inplace(element97& res, const element97& x) noexcept {
  res.value = (res.value + x.value) % 97u;
}

//--------------------------------------------------------------------------------------------------
// mark
//--------------------------------------------------------------------------------------------------
inline void mark(element97& res) noexcept {
  res.marked = true;
}

//--------------------------------------------------------------------------------------------------
// is_marked
//--------------------------------------------------------------------------------------------------
inline bool is_marked(const element97& e) noexcept {
  return e.marked;
}
} // namespace sxt::bascrv

// main
using namespace sxt;
using namespace sxt::mtxbk;

int main() {
  using E = bascrv::element97;
  E* bucket_sums = nullptr;
  const E* generators = nullptr;
  const uint8_t* scalars = nullptr;

  // f(bucket_sums, generators, scalars, 0); // uncomment this line and things compile. ¯\_(ツ)_/

  kernel<<<1, 1>>>(bucket_sums, generators, scalars, 0);

  return 0;
}
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcWU1v8j4S_zTmMgIlDhA4cIBStL33rz0ixxnAW8dm_dLSy_PZV3aSEmha8Uj_Z-luhCBkxvPun8cOs1bsFeKCTFZksh4w7w7aLIwqvFGDQpfvi2f2guAOCHhi1VEicF0iHLV1WEKBUr-NSLImybL-ftrBE3BdHYVEQh_gCfbowB2EBTRGm4Z3mjSf-m_2CFwytSd0RegKhifgvmQw5DAchrvh_uiHzPADyda22uYJDK0rSbbm9QiaQMWEGnFeC-wII9kS3phRQu3D7cNf6yW8orFCKxAWFL6hAXdgKjopmUPr4MiME0zKd7D-eNQm-NoOSukoBTJZDf_p1YvSb6q2sCGTybo2wWjtyDjJGMvy6Szd0aIk2ZLQjTWc0A3XiuPRbQu_JzT7v_O-tSdbpumUZMssaK8LIFuC0lAxxw9C7WHnFXdB8k4b4ExKcBoIzV_QKJSE5rVAgDSdAskfAKAhZQ_1Jw11lpLssfnQWeH5C7qt9ZUNtD0qNMxpE_9ZziSrbxNC5yRbtRrCVWsgk8df4bp2hgZf0rz2wWH45UyVomQOwWF1DBEEsVfaYBmpWllnmFDOhhFgmRN2J7AMMXwT7gDPQLI1PH5EDgBo7eh2u5e6YHK7hVctytZtOnsmdAlXPkZFECmX7tYEL5SbbV0gn_3_7PflRSaPnwIwiSntBqBAzrzFkDN7cqHGs2XBLDev9T1KrFC5eU5oDqXGbiDew6iGoZNqoJNo0EdESfbQIxKeQ7q_dKLH_DRra7HH-lL7QuL2w5yZwRhBDFVCc3jTXgbIA6FemRTlrZV8LffsZpo1Jn-r-osCrR-m9axqk44lvB1QNQgcrdIGImiM-qE3fq-8iwDQoBK0AbaH1ueQr1C3I3hSsGPcBfNEAHuvuK4id4QQobBXEQChG0I3sPv9CdoO7aoSNuoCpsrwT-1tu-yMgDxQstyQycOW0Bl5zMgsI7MxofNtENRnXe18GM5iCt-ElGA954jlxfr2DzQYkDMMaDFR72r8hidC8wq8DWjbp-UC5T8wM5L-KrxyvpHzAbazUTJKIDjRAD3NknmaJWOaZimhq122K_l8mrNZXg7TX3g6pr86THQyHqXpjIYoRi3PzOwxzFE4zabb6Xh45EMplD8N98o3LAeDrIRKlwFil2GlF6dmbVfWMSmxXIsI44RuvDWEbgqhvimuxnddhgnXz0czobj0ZZzpQltnkFUfk7tlijXQlObo0DeUW1cK5c4jL6l1ddsrwYpVaI-MI1zhF5B81dU9_ONXn6f3sKALvFwyaztQe40RYfky-G8vImaFxWkKDXx9LErTiGTneN6MeKwsL2ifGBTuvxu8FeooGcczT_zpcpF8FbwLaRclKifce5xwwdw1KI2n6O4wNKmx7QqcllW4ZZZkDzEwZ2kVMy-tsmstwm4DGYNLeLOCQmvZ1UHydfe-xcYvq7hnDtWdRDuFekZW7lS8nNN1-zr8X-1ZPl1exe1MCRLV3h2uSo7QWbCJ0HnXjqyH3rGmj9xa00drNHeT1Yn_7YGM0dv9jYH70cH5toxjMfatBPWGdPt7K8L_AuTP83vYYJ3x3J1t6JZHKKmMbh28Mumxk9yATlCjWlwJdkzaM73lso45wTuir5D2AwQ7KgEMOm9UeJZ0QS-A6fpSfrRCH-u6JNk6fOisng2d_cc06KqfnjVmayhxx7x0_RB7v2q4WiPvYIhQsceOYPRpxe4G9nLFv6DgVwk2aEexnGIOgsD2L11Bcx9Rik5gnvuvMPUuqQltx53zUXc-f3sS5rmH4U8Pf2gL7xz-ujO9MfynL56_35yWU2duvP_45Hy03fdPUmcHcHuybkxK5wFdwelnpyXuTO6cj2Z39DkRXwe801s402k9fkJIzxu6-8U1tj7dneXvoX_ssLCJ8u925l9uMCvWngjFw7DrsatvaOeO_7IDE8rVYtu9-bk3jWIeY4n0HzyfO8fHq91UHKS8lEdnOmx1DB8vN1jf8n7ec_Wyn_dEP-kstDHqz7xXuaq1pKfI2tPAQbnIynk2ZwNcpNP5JM_HaTYfHBZpOWXFJE-zclxMxnlRpAWdJ-Wc8SQvUpoNxKI-9qR5klCazkcsmTFeTGfIJywdzydknGDFhBxJ-VqNtNkPhLUeF9N8kuQDyQqUNr4FpbQ5paVksh6YReAfFn5vyTiRwjp7luCEk7h4iEe28Y1d-wagfsVxYKqMR_DtuSNwbQxyJ99HA2_k4uDc0dYvBQnd7IU7-GLEdUXoJuhofoZHo_-F3BG6iSZbQjfR6v8EAAD__4QT-F4">