<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/63199>63199</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
clang CUDA. error: no matching constructor for initialization of std::shared_ptr
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
burlen
</td>
</tr>
</table>
<pre>
It seems that clang cuda wrappers overload `free` from the C stdlib in a way that causes some valid C++ code to fail when compiling for cuda. The issue specifically is with overload resolution in the `std::shared_ptr` constructor when using `free` as the deleter.
Here is small reproducer:
```C++
#include <cstdlib>
// note: the noinline issue was fixed in clang may '23
#undef __noinline__
#include <memory>
#define __noinline__ __attribute__((noinline))
int main(int, char**)
{
unsigned int *tmp = (unsigned int*)malloc(16*sizeof(unsigned int));
// fails with -x cuda
auto ptr = std::shared_ptr<unsigned int>(tmp, free);
return 0;
}
```
To reproduce compile this with:
```
clang++ -std=c++17 -x cuda --cuda-gpu-arch=sm_75 test.cpp
```
The errors are:
```
test.cpp:14:16: error: no matching constructor for initialization of 'std::shared_ptr<unsigned int>'
auto ptr = std::shared_ptr<unsigned int>(tmp, free);
^ ~~~~~~~~~
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:535:7: note: candidate constructor not viable: no known conversion from 'unsigned int *' to 'const weak_ptr<unsigned int>' for 1st argument
shared_ptr(const weak_ptr<_Tp>& __r, std::nothrow_t) noexcept
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:231:2: note: candidate template ignored: couldn't infer template argument '_Deleter'
shared_ptr(_Yp* __p, _Deleter __d)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:248:2: note: candidate template ignored: couldn't infer template argument '_Deleter'
shared_ptr(nullptr_t __p, _Deleter __d)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:311:2: note: candidate template ignored: could not match 'shared_ptr<_Yp>' against 'unsigned int *'
shared_ptr(const shared_ptr<_Yp>& __r, element_type* __p) noexcept
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:463:2: note: candidate template ignored: could not match '_Sp_alloc_shared_tag<_Alloc>' against 'unsigned int *'
shared_ptr(_Sp_alloc_shared_tag<_Alloc> __tag, _Args&&... __args)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:204:7: note: candidate constructor not viable: requires 1 argument, but 2 were provided
shared_ptr(const shared_ptr&) noexcept = default; ///< Copy constructor
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:214:2: note: candidate constructor template not viable: requires single argument '__p', but 2 arguments were provided
shared_ptr(_Yp* __p) : __shared_ptr<_Tp>(__p) { }
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:351:2: note: candidate constructor template not viable: requires single argument '__r', but 2 arguments were provided
shared_ptr(const shared_ptr<_Yp>& __r) noexcept
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:359:7: note: candidate constructor not viable: requires single argument '__r', but 2 arguments were provided
shared_ptr(shared_ptr&& __r) noexcept
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:368:2: note: candidate constructor template not viable: requires single argument '__r', but 2 arguments were provided
shared_ptr(shared_ptr<_Yp>&& __r) noexcept
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:380:11: note: candidate constructor template not viable: requires single argument '__r', but 2 arguments were provided
explicit shared_ptr(const weak_ptr<_Yp>& __r)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:387:2: note: candidate constructor template not viable: requires single argument '__r', but 2 arguments were provided
shared_ptr(auto_ptr<_Yp>&& __r);
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:395:2: note: candidate constructor template not viable: requires single argument '__r', but 2 arguments were provided
shared_ptr(unique_ptr<_Yp, _Del>&& __r)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:412:17: note: candidate constructor not viable: requires 1 argument, but 2 were provided
constexpr shared_ptr(nullptr_t) noexcept : shared_ptr() { }
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:268:2: note: candidate constructor template not viable: requires 3 arguments, but 2 were provided
shared_ptr(_Yp* __p, _Deleter __d, _Alloc __a)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:287:2: note: candidate constructor template not viable: requires 3 arguments, but 2 were provided
shared_ptr(nullptr_t __p, _Deleter __d, _Alloc __a)
^
/usr/lib/gcc/x86_64-redhat-linux/12/../../../../include/c++/12/bits/shared_ptr.h:202:17: note: candidate constructor not viable: requires 0 arguments, but 2 were provided
constexpr shared_ptr() noexcept : __shared_ptr<_Tp>() { }
^
```
After a bit of digging I found that clang cuda wrappers overload free for use on the device. For instance:
``` C++
#include <cstdlib>
int main(int, char**)
{
// this line give some additional info. it seems the CUDA
// wrappers overload free and type of free cannot be deduced.
using free_t = decltype(free);
return 0;
}
```
errors out with:
```
test.cpp:18:29: error: reference to overloaded function could not be resolved; did you mean to call it?
using free_t = decltype(free);
^~~~
/home/bloring/work/llvm/llvm-17_5_25-install/lib/clang/17/include/__clang_cuda_runtime_wrapper.h:405:17: note: possible target for call
__device__ void free(void *) __attribute((nothrow));
^
/usr/include/stdlib.h:568:13: note: possible target for call
extern void free (void *__ptr) __THROW;
```
Showing the device overload. Adding an explicit cast on free can be used to help the compiler find the right overload, `auto ptr = std::shared_ptr<unsigned int>(tmp, (void (*)(void *) noexcept(true))free);`
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzMWk1v47oV_TXM5iKGTdmSsvDCScZ4b1Wgb4qiK4IWryR2KFKPpOLkLfrbC1KyJSdxZjKZNjYMxbZE8vLwnMuPE-6crDTimqxuyer-ine-Nna966xCfbUz4mn9uweH2DjwNfdQKK4rKDrBYW9526J1YB7QKsMFkHReWkSSzqG0pgFfI9yB80LJHUgNHPb8aaiHdw4dONMgPHAlBdwRekvoLRRGIHgDJZcK9jVqKEzTSiV1BaWxse0ZfK0RpHMdgmuxkKUsuFJPIB3spa_HkCw6ozovjQ4BhIBIOndekGRDko2ruUXBWm9DzIXRztuu8Mb2DXcuNDrpFXexCoEKPdoZmd-T-aa__oY2RASu4UqBxdYa0RVoQ0P9Y-m8fw8dHX6lidSF6gQCSe6KHiuSfJlWTeiW0C1o45EkmxiBNlIrqQ8Y7LmDUj6iCJ3sh6jhT0BoRpNjO50WWAJjh7KMvRZCg42xT2MENBFYhpamBYEx7r2Vu86HamhOaH64S-hNeE_il9pDw6UmNJfaE3oHRc0toZv4PjyaDYgAAHQ60jJ0xwOhG9-0QJJ7IDSf3urLB8BNQWi-SAndOPkXmvLFgzGm5HYaVmhogDZwbWDO9WNk2PgI77yB1tsYwGvMSe5Omkq-EJr7pg39jLx5tWGLvrMa5uOt7P4ZT6ZFvpqRUoMeEHw90P0lx_qvkQmDrK5j6PdF_3WRHToK19fhz3XVdtfcFjVJ7l3DshV4dH5WtO1bUdUIaK2xDrjFc1EcK0o2i2W4pIHFsVz4oA003Bd1kNpUgUHrUksvuZJ_8ahgUwZK_-AYZP-bIQSy-gL_ObyOAu2cJXQbtEu3VVEQun3MU5Yury2KmvtrJXX3SOh2QQndzmYvL4MECd0OQ3R4eCe9I3Q7RjsL471KViTZZD2AfV4ouBZScI8nMGrj4UHyncIB7G_a7ENO1Q9oXUA1pmpCs-eiIzQLeZjQLNYHe-TfzoIdh2vhPHBbdQ1qP6A1AZnmL-phX9tYPAXGbED7ODja-NqaPQvSBW3wscDWjyPwabjTZBGur-PusWlV-CArbSyKeNd0SmhCMw9Sl2jHhw5IBYTZfT-lnNA2vE7wY_9qCd0AY5GZhzLAmDim0c8FZ5l_Hji6U6r1lvnLxSdZ_Ax5ooRjiozZb5q1AiF6-fGKyyCu13X8Bmy9Jl-t9ahKVBjGgvmnFo8EvChZLtPkw8iyP1oWlxJsqNzzKqCxicuLd8B8qtnv1Aos_hL4urGVIzQlNJ3NZmGJFb9fAnHpfPkzs43FPztp0cFinBboHew6DxT2YcHcWvMgBYopQd-mZ8BnpF6c0QWWvFOeJLfDei68kzu4M-3TNLpLADKugM7wdArkkbNnEA07E_UsTQZVZiPAh3vuPNRvzC83ENpj7DQzDPN1fngku4XjwvVzU-vqjdT6YWDtB4H9fpK9qGyarG4-JPdfg-EJgs8ywEWilr6x-vl0Cr5OvgtFMp-HfeLicqDEx1bJQvrv7WeeifoisMwumJVhZ36ekyeb7s9D8GZ1wQh2Wv7Z4QTDYefzAsvPB3K5CCAu_k_LyFgfPrb2zDbx2TLy2bL9slY39FdNLcnIsB-A8B3HD2H7EnY0YdtyGWyjvyrxfQiz75xKXCBs8w-KdP4uuM6o9IU4z21EXur02WvE8vRUelOGQeCwkx5MCUJWldQV_A6l6bT4AbOrtIjx1LNzCEYPxtCDLHAG23h27TzXxeRY_DQMeLcP9H4fZbA3olEQ7aJKPmBvunEhpJdGcwVSl2YGcvT5EO7-cf_SJTmDAQ9wPbUYYIw_FFwHbuwCHqIrUMwmzk601MJj7LB5L1R_rpS_zy0JcZ23JgZbwnT-TYdk6k3EFHtz4k1YLNGiLqIheeg1Cig7XURHYjxG2mHvNT6gIMktCCngyXTQINehcMGVAulJsv05LAKPp35DbZog6Z0yVuqK0O3e2G8hUaiHZvhzvcjYitHVdWSiUsc0MthC20V2kh4YizdY4DuznfayQTaMeT9_z1cvUkNrnJM7heC5rdD3Dm1oLAbKWK8IxuDBSDGYKXn83JN26iQefMR48v_Msjsv60OKHDvSi6e3SeLEuUh-PGR89Gj1GC5M4mV9dgpBf_3t73_750jHU1b9UZt9GNoxJRzJM4ONEOEe1-OmouDOQ3RievUEMnUORSBOjaqNFQ2un4VSxvyEYGVV-2PNIR-QdP4Rr2vsaj6klNOxOm4Rae5tNzi9E6am8yuxTsRNcsOvcL1I85Qu02WaX9XrcrWiPKe7RGRpiUm5yFYrjongK1ruBMcruaZzmszTeb7Iltkyn-1ElhWCijy_wYTnSJZzbLhUs8DtmbHVVfS-12myuLm5UnyHysV_Y6BU4743xgmlZHV_ZddRD7uucmQ5V9J5N9bipVe47lN9SHuzn7UmXwH7qrNqXXvfunAnZtFK-rrbzQrTPFNra82_sfCBySH0MDXHrv03AAD__8dbU3w">