<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/100802>100802</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[libc++] std::memcpy does not work in HIP __device__ code
</td>
</tr>
<tr>
<th>Labels</th>
<td>
libc++
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
AngryLoki
</td>
</tr>
</table>
<pre>
This code (seen in pytorch) succeeds with stdlibc++, but fails with libc++:
```c++
#include <cstring>
__device__ void test1( void* dest, const void* src, std::size_t count ) {
// this fails
std::memcpy( dest, src, count );
// but this works
memcpy( dest, src, count );
}
```
Fails with error: reference to __host__ function 'memcpy' in __device__ function, see https://godbolt.org/z/h5nEnbb68
The issue lies between these lines:
https://github.com/llvm/llvm-project/blob/c80c09f3e380a0a2b00b36bebf72f43271a564c1/clang/lib/Headers/__clang_hip_runtime_wrapper.h#L142-L145
For math functions [__clang_cuda_math_forward_declares.h](https://github.com/llvm/llvm-project/blob/c80c09f3e380a0a2b00b36bebf72f43271a564c1/clang/lib/Headers/__clang_cuda_math_forward_declares.h) exists and adds definitions for all math functions. But for stdlib functions there is no such file.
What happens in stdlibc++:
```c++
// Part1: /usr/include/string.h
extern void *memcpy (void *__restrict __dest, const void *__restrict __src,
size_t __n) noexcept (true) __attribute__ ((__nonnull__ (1, 2)));
// Part 2: "/opt/compiler-explorer/libs/rocm/6.1.2/include/hip/amd_detail/amd_device_functions.h"
static inline __attribute__((device)) void* memcpy(void* dst, const void* src, size_t size) {
return __hip_hc_memcpy(dst, src, size);
}
// Part 3: /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/cstring
namespace std {
using ::memchr;
using ::memcmp;
using ::memcpy;
...
```
In libc++ part 2 and part 3 are swapped, because `#include <algorithm>` in libc++ includes `<.../c++/v1/cstring>`, which results in
```c++
// /opt/compiler-explorer/clang-rocm-6.1.2/bin/../include/c++/v1/cstring
namespace std { inline namespace __1 {
using ::size_t __attribute__((__using_if_exists__));
using ::memcpy __attribute__((__using_if_exists__));
using ::memmove __attribute__((__using_if_exists__));
...
// /opt/compiler-explorer/libs/rocm/6.1.2/include/hip/amd_detail/amd_device_functions.h
static inline __attribute__((device)) void* memcpy(void* dst, const void* src, size_t size) {
return __hip_hc_memcpy(dst, src, size);
}
```
Adding `extern __attribute__((device)) void* memcpy(void* dst, const void* src, size_t size);` before `<algorithm>` in `__clang_hip_runtime_wrapper.h` solves the issue.
In general it looks like a new file like `__clang_cuda_stdlib_forward_declares.h` with memcpy and memset could solve the issue without breaking anything.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzcV81u27wSfRp6M4ggUbbsLLxImgot0EUXBe6SoKixxRuaFEjKTvr0F0P5R3aCtAWKiw8fYMQxf4Yzh3PODGUIemsR12zxyBZPMznEzvn1g93612_uWc8a176uf3Q6gHItAuOrgGhBW-hfo_OqY_wewqAUYhvgoGMHIbZGN4rxx_T5BM0QYSO1Oc5PJssHlj-x_PS3ysfPaXoc5aW2ygx0evlJhei13bLy83SnEC3utUIhYO90CxFDLBhfpV-MP0CLIZIrytkQz6PBKxoMsSVPyoegf6KIoNxgI1BgbHl0AgCA8ZrxGiKBkcI5Tp2373Cn-lc69nTc8YCzQVZO7d1aJpyS9YPzzyfrf2CULZ9ugJxiVF9uAL13npUP4HGDHq1CiA6E6FyIQsBmsCpqZ4Hx5en4JV35BObTmuQQInQx9oFgSKFsXds4EzPnt4zXPxmvu4X9bJumWk1d-tEh6BAGBKMxQIPxQMkVOww0ZDHcZMjNKTp2Q5Mpt2O8NmZ_-rrrvfsvqsh43RjXMF6rVa7y-02J5SqXueRNnjdl1WCzWfLNvOTLQi6quSpoqZGWnDaaNn5B2aIPjNdCpBnR6V74wUa9Q3Hwsu_RZx3j5bdizu--FfPFFebOw07G7gxXALZ4PJlSQysFTYuN8wfpW9GiMtJjyDq2eGJ89Q-I90Mn-T3giw4xgLQtyLYN0OJGWz3GunEepDE3EGTwSILg_FEpJuDEDj2lBFhHotLBRhvMpoj-p5MROoLdBkrJa7Epf6EkiWbfpY8FZT_j9RA84_VRXxivR3XJunEDvkT0dpQUxo_8Jg08jQjhkbaomLhxqzFvlozUvRAfAI6aI4QlNK3DF4U98XoV_YA0JoSM0etmiEQ8xleMr4SwztrBmHGkoHM5acH4OSvCbdzAx8A547XrKWOU2_XaoL_Dl944j37MBUoB7xSlWJUVGb9CqdM947XcUS5Eqc35RxKHy0V3dFByIUQZtQJtidXXEY0BjXtH98_6fNa-s4x_pOIjkPT1Rro9xsGTfhF7OyXOhtsrST3ufUdQ3-BYXiXQSJ6tUozXL6tKVPM7o-3wcre1A-N1QUTLsrd_LpCey-W4-FTm0rlW7jD0UiFl-zSwIWi7hUvx6fykwNxO7voPJvvX82SWZR9Uka92Ur-hTymVyJ_-LUF6hHAgfrap9KOSQ0AgK1d1XJqt8zp2O6rkVU5Untg9LgxpX_kpS2BdINpPIUr76axDp1UHHsNgYtKG39CCD2mQBPCOWHB34kCj7Ud3t__13Z04cJkQorhc6vRiztLwhi1CpHVCb8SovzRxxfy3F_xX7Ozc_l36_trQJa1-G_-_J0P_Cg16j44PbZvup8qPter_EBa5VuXQ4MZ5PDL0HTqzKv-4Z6pyCM7sMZX9sQ_MbqRmixa9NKAjGOeeAxj9jCDB4iG1BuPvyUmpWxl7gvf6lSofG-AjI0i3drgLmHp-047-XNxJi90QofEonwloaV9jRx3CrF2X7X15L2e4Lpac57yolvNZt24Q5ZLjQuGqaDctFjyXebVZ8CqXWOFiptc85_N8yavinhfFPKvyKr_n2JabfNmUfMHmOe6kNhn1dtRCz5Iz6yLPVzmfGdmgCenBxvn0ocXpAefXqSNshm1g89wQFS92oo4mPfUm2xZPtw8YaB1SBxbTO4Su8svX79Pen16Cs8Gb9R-3pykQ4vQxlv2a_y8AAP__QL18sw">