Disable [[no_unique_address]] for clang and mdspan

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

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 <cuda/std/cassert>
+#include <cuda/std/mdspan>
+
+#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<empty, int>;
+  _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;
+}