diff --git a/libcudacxx/include/cuda/std/__cccl/attributes.h b/libcudacxx/include/cuda/std/__cccl/attributes.h index 192dcaae165..4352b3d1590 100644 --- a/libcudacxx/include/cuda/std/__cccl/attributes.h +++ b/libcudacxx/include/cuda/std/__cccl/attributes.h @@ -67,6 +67,12 @@ # define _CCCL_NO_UNIQUE_ADDRESS #endif +// Passing objects with nested [[no_unique_address]] to kernels leads to data corruption +// This happens up to clang18 +#if !defined(_CCCL_HAS_NO_ATTRIBUTE_NO_UNIQUE_ADDRESS) && defined(_CCCL_COMPILER_CLANG) +# define _CCCL_HAS_NO_ATTRIBUTE_NO_UNIQUE_ADDRESS +#endif // !_CCCL_HAS_NO_ATTRIBUTE_NO_UNIQUE_ADDRESS && _CCCL_COMPILER_CLANG + #if _CCCL_HAS_CPP_ATTRIBUTE(nodiscard) || (defined(_CCCL_COMPILER_MSVC) && _CCCL_STD_VER >= 2017) # define _CCCL_NODISCARD [[nodiscard]] #else // ^^^ has nodiscard ^^^ / vvv no nodiscard vvv diff --git a/libcudacxx/test/libcudacxx/libcxx/containers/compressed_pair.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/containers/compressed_pair.pass.cpp new file mode 100644 index 00000000000..90bdeec7d15 --- /dev/null +++ b/libcudacxx/test/libcudacxx/libcxx/containers/compressed_pair.pass.cpp @@ -0,0 +1,45 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++11 +// UNSUPPORTED: nvrtc +// UNSUPPORTED: msvc && c++14 +// UNSUPPORTED: msvc && c++17 + +#include +#include + +#include "test_macros.h" + +// We are experiencing data corruption on clang when passing a mdspan mapping around where on of the subtypes is empty +struct empty +{}; +struct mapping +{ + using __member_pair_t = _CUDA_VSTD::__detail::__compressed_pair; + _CCCL_NO_UNIQUE_ADDRESS __member_pair_t __members; +}; + +__global__ void kernel(mapping arg1, mapping arg2) +{ + assert(arg1.__members.__second() == arg2.__members.__second()); +} + +void test() +{ + mapping strided{{empty{}, 1}}; + kernel<<<1, 1>>>(strided, strided); + cudaDeviceSynchronize(); +} + +int main(int, char**) +{ + NV_IF_TARGET(NV_IS_HOST, test();) + return 0; +}