Fix `<chrono>` and `<atomic>` build errors with clang-cuda. #304

wmaxey · 2022-08-19T04:03:24Z

This fixes compilation in clang-cuda when atomics are included.

EDIT:
Testing support has been split into a separate branch and will include Docker changes to add support for testing.

jrhemstad · 2022-08-19T11:55:19Z

include/cuda/std/detail/libcxx/include/__config

@@ -1662,7 +1662,7 @@ extern "C" _LIBCUDACXX_FUNC_VIS void __sanitizer_annotate_contiguous_container(
 #endif

 // CUDA Atomics supersede host atomics in order to insert the host/device dispatch layer
-#if defined(_LIBCUDACXX_COMPILER_NVCC) || defined(_LIBCUDACXX_COMPILER_NVRTC) || defined(_LIBCUDACXX_COMPILER_PGI)
+#if defined(_LIBCUDACXX_COMPILER_NVCC) || defined(_LIBCUDACXX_COMPILER_NVRTC) || defined(_LIBCUDACXX_COMPILER_PGI) || defined(__CUDACC__)


Should we add a _LIBCUDACXX_COMPILER_CLANG_CUDA definition that we can reuse elsewhere?

I don't see why not. I'll edit that in.

jrhemstad · 2022-08-19T11:55:41Z

.upstream-tests/test/support/cuda_space_selector.h

+# define SHARED
+#endif
+
+#if defined(__clang__) && defined(__CUDA__)


e.g., https://github.com/NVIDIA/libcudacxx/pull/304/files#r950113474

We generally try to refrain from using libcudacxx macros in the tests. This allows us to ensure the behavior matches up.

robertmaynard · 2022-08-19T17:25:56Z

.upstream-tests/test/CMakeLists.txt

+
+  set(LIBCUDACXX_TEST_LINKER_FLAGS
+    "${LIBCUDACXX_TEST_LINKER_FLAGS} \
+    -L${CUDA_TOOLKIT_ROOT_DIR}/lib64 -lcuda -lcudart")


You should use:

find_package(CUDAToolkit) target_link_libraries( tests PRIVATE CUDA::cuda_driver CUDA::cudart)

You're a saint Robert.

I thought this was done for me when enabling CUDA as a language. That's not the case?

If CMake is generating the build files it will make sure that cudart is part of the link line ( with controls over static / shared ). It seems that in your case you are passing this information down to lit which does the build generation so you might need to stick with the 'bad' ( original ) approach.

robertmaynard · 2022-08-19T17:26:45Z

.upstream-tests/test/CMakeLists.txt

@@ -42,6 +42,18 @@ set(LIBCUDACXX_TEST_COMPILER_FLAGS
  -I${CMAKE_SOURCE_DIR}/include"
  CACHE INTERNAL "Flags for libcxx testing." FORCE)

+if (${CMAKE_CUDA_COMPILER_ID} STREQUAL "Clang")
+  set(LIBCUDACXX_TEST_COMPILER_FLAGS


string(APPEND is the preferred CMake style

griwes · 2022-09-27T18:28:45Z

.upstream-tests/test/support/cuda_space_selector.h

@@ -101,7 +105,7 @@ struct device_shared_memory_provider {

    __device__
    T * get() {
-        __shared__ alignas(T) char buffer[shared_offset];
+        alignas(T) __shared__ char buffer[shared_offset];


This being needed is really funny to me.

griwes · 2022-09-27T18:29:27Z

.upstream-tests/utils/libcudacxx/compiler.py

@@ -146,6 +146,9 @@ def _initTypeAndVersion(self):
        if self.type == 'nvcc':
            # Treat C++ as CUDA when the compiler is NVCC.
            self.source_lang = 'cu'
+        elif self.type == 'clang':
+            # Treate C++ as clang-cuda when the compiler is Clang.


typo ("treate")

Also "treat c++ as clang-cuda" doesn't particularly make sense, did you mean to say "as CUDA" as above?

griwes · 2022-09-27T18:31:23Z

.upstream-tests/utils/libcudacxx/test/config.py

+        # I don't think this is required, since removing it helps clang-cuda compile and libcudacxx only supports building in CUDA modes?
+        # if self.cxx.type != 'nvcc' and self.cxx.type != 'pgi':
+        #    self.cxx.compile_flags += ['-nostdinc++']


Yeah I think you are right. Change the comment to state this as a fact, with a note to reenable it if we ever decide to start actually testing in a C++ mode.

griwes · 2022-09-27T18:32:02Z

.upstream-tests/utils/libcudacxx/test/config.py

-                self.cxx.link_flags += [abs_path]
-            else:
-                self.cxx.link_flags += ['-lc++']
+        # Device code does not have binary components, don't link libc++


Same as above, add a note to reenable this if we decide to start testing in a C++ mode.

griwes · 2022-09-27T18:32:49Z

include/cuda/std/chrono

@@ -50,7 +50,7 @@ system_clock::time_point system_clock::now() _NOEXCEPT
 {
 #ifdef __CUDA_ARCH__
    uint64_t __time;
-    asm volatile("mov.u64 %0, %globaltimer;":"=l"(__time)::);
+    asm volatile("mov.u64 %0, %%globaltimer;":"=l"(__time)::);


I am really disappointed that NVCC accepts the old inline asm.

wmaxey · 2022-10-11T20:20:49Z

I've split the fixes from the CMake changes. I'll add those into another branch and will add them into some Dockerfile steps later for proper testing support. I haven't done this because they have the side effect of breaking our containers, I'm just merging the fixes instead.

wmaxey requested review from griwes and miscco August 19, 2022 04:03

wmaxey self-assigned this Aug 19, 2022

wmaxey added this to the 1.9.0 milestone Aug 19, 2022

jrhemstad requested a review from robertmaynard August 19, 2022 11:51

jrhemstad reviewed Aug 19, 2022

View reviewed changes

robertmaynard reviewed Aug 19, 2022

View reviewed changes

griwes approved these changes Sep 27, 2022

View reviewed changes

wmaxey force-pushed the bugfix/clang_cuda_atomic_support branch from 525abbe to 676511d Compare October 11, 2022 20:05

Fix issues preventing compilation in atomic headers with clang-cuda

009be41

wmaxey force-pushed the bugfix/clang_cuda_atomic_support branch from 676511d to 009be41 Compare October 11, 2022 20:15

wmaxey changed the title ~~Add support for testing libcudacxx with clang-cuda enabled.~~ Fix <chrono> and <atomic> build errors with clang-cuda. Oct 11, 2022

wmaxey merged commit 2253815 into main Oct 11, 2022

wmaxey deleted the bugfix/clang_cuda_atomic_support branch October 11, 2022 22:47

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix `<chrono>` and `<atomic>` build errors with clang-cuda. #304

Fix `<chrono>` and `<atomic>` build errors with clang-cuda. #304

wmaxey commented Aug 19, 2022 •

edited

Loading

jrhemstad Aug 19, 2022

wmaxey Aug 19, 2022

wmaxey Aug 19, 2022

jrhemstad Aug 19, 2022

wmaxey Aug 19, 2022

robertmaynard Aug 19, 2022

wmaxey Aug 19, 2022

robertmaynard Aug 22, 2022

robertmaynard Aug 19, 2022

griwes Sep 27, 2022

griwes Sep 27, 2022

griwes Sep 27, 2022

griwes Sep 27, 2022

griwes Sep 27, 2022

wmaxey commented Oct 11, 2022

Fix <chrono> and <atomic> build errors with clang-cuda. #304

Fix <chrono> and <atomic> build errors with clang-cuda. #304

Conversation

wmaxey commented Aug 19, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

wmaxey commented Oct 11, 2022

Fix `<chrono>` and `<atomic>` build errors with clang-cuda. #304

Fix `<chrono>` and `<atomic>` build errors with clang-cuda. #304

wmaxey commented Aug 19, 2022 •

edited

Loading