[FEA] make raft sources compilable with clang #424

MatthiasKohl · 2021-12-14T22:42:47Z

This makes RAFT sources compilable with clang.
It fixes some fragile code (using static const instead of static constexpr or %laneid in PTX relying on quirks in nvcc which make this happen).

RAFT is still not compilable with clang entirely though due to the dependencies:

cub has this issue before 1.14: Fix clang compile error NVIDIA/cub#335
libcudacxx has issues with atomic, which should be fixed in >= 1.7.0-ea (wasn't able to verify this yet)
libcudacxx has issues with variadic CUDA functions, which is apparently fixed by passing -Xclang -fcuda-allow-variadic-functions to clang (wasn't able to verify this yet)
cooperative_groups from CUDA does not work with clang 11.0 / 11.1 but only with >= 13

EDIT: this is necessary to close #84

… cg, libcudacxx)

MatthiasKohl · 2021-12-14T22:44:42Z

question to reviewers: what would be the easiest way to bump up versions of tools and dependencies ?
In particular, to make this really happen, we would need thrust/cub >= 1.14, libcudacxx >= 1.7.0-ea and clang-tools >= 13.

With the Rapids cpm system, this seems prohibitively hard. Why?

cjnolet · 2021-12-20T22:08:02Z

@MatthiasKohl, I know you've put this down in the meantime due to some issues w/ dependent versions in RAPIDS. I still see value in merging the c++ code changes and keeping the script in a PR until clang support is ready to go. What do you think?

MatthiasKohl · 2021-12-21T17:15:02Z

Yes I think it makes sense to add these changes to RAFT. But it's hard to test because of these issues in dependencies meaning that for many files, clang doesn't even get to the RAFT part. In fact after I upgraded to clang 13 and updated some of the dependent packages, I still discovered new issues within RAFT sources because they were simply not compiled before.
So what we could do is to just ignore the errors for now and gradually change things as more and more code will be covered with newer packages and newer clang versions

…udacxx/cuco)

MatthiasKohl · 2022-01-03T16:08:50Z

@cjnolet I've added a few changes I had in the pipeline which make things compatible with CUDA 11.5, plus a few other RAFT sources which can only be compiled with a newer clang version and updated thrust.

Can you review these changes so that we can already merge those?

Right now, libcudacxx and cuco cannot be compiled device-side.
I can still check whether they can be compiled host-side which would already be nice.
But I feel like that isn't really enough for a project like RAFT where probably 50% of code is device-side, and especially the important code is device-side, so before spending more time on this, I'd prefer having the dependent projects work with device-side compilation.

MatthiasKohl · 2022-01-03T16:08:57Z

rerun tests

cjnolet

Since RAFT is still somewhat coupled to cuml's C++ code, can you open a PR in cuml w/ the get_raft.cmake pinned to this branch just so we can make sure it doesn't break? Overall the changes look good.

cpp/include/raft/spectral/kmeans.hpp

cpp/scripts/run-clang-compile.py

cpp/include/raft/cuda_utils.cuh

MatthiasKohl · 2022-02-01T15:31:35Z

@cjnolet sorry I didn't see your comments on this PR (somehow missed the notifications), I'll address them ASAP

…copyright

…g-compilable

MatthiasKohl · 2022-02-14T13:52:02Z

@cjnolet you might have seen this draft PR on cuml, which I'm using to make sure that cuml CI is still passing with these changes: rapidsai/cuml#4571

Note that I had to add some more changes to raft in order to make sure that RAFT sources themselves are compiling with Clang: - cpp/include/raft/cuda_utils.cuh lines 697-728 : I added overloads for the __ldcv/__ldcg/__stwt CUDA functions which are unknown to clang

https://github.com/rapidsai/raft/pull/424/files#diff-ca87b707f7d1c64a1085887c53fc4e80799a0b82ccc19a4aff3471141ad79dbb lines 317-320 : I excluded the __shared__ alignas construct from clang compilation since Clang cannot compile that unfortunately.

MatthiasKohl · 2022-02-14T18:43:02Z

@cjnolet CI on cuml is passing with the latest changes from this PR (see rapidsai/cuml#4571)
It seems like CI here on the RAFT PR isn't being run for some reason. EDIT: actually, it is being run but unfortunately, it looks like one of the Rng tests (MeanError) is not passing.

Apart from that, I'm still awaiting your final review for the latest changes since the compiler-specific stuff might not be what you'd like to do for RAFT in general.

cjnolet · 2022-02-14T18:52:23Z

rerun tests

cjnolet

I'm still thinking through the best way to isolate and document the files that can be safely compiled with a CUDA vs non-CUDA-enabled compiler while keeping the API consistent.

Thanks again for these changes! I really just have a couple minor-ish questions and one small change.

cjnolet · 2022-02-14T19:15:11Z

cpp/include/raft/cluster/detail/kmeans.cuh

+ * Ceiling division operation
+ * Re-declared here from raft/cuda_utils.cuh since we cannot include CUDA files
+ */
+template <typename IntType>


since we cannot include CUDA files

Can you explain what you mean by this? I think I'm misunderstanding only because I notice this file has kernels in it. Is it the __CUDA_ARCH__ that's not allowed?

Thank you for raising this again!
Given the comment, I must have thought that this file was compiled from a .cpp file (thus with a non-CUDA-enabled compiler). I actually don't think this was the case, and these changes don't make any sense to me anymore, so I reverted them and they don't seem to affect clang compilation at all.

cpp/include/raft/cuda_utils.cuh

cpp/include/raft/matrix/detail/linewise_op.cuh

cpp/include/raft/cuda_utils.cuh

…to use of ld/st CUDA intrinsics with cache hints

MatthiasKohl · 2022-02-16T13:59:50Z

@cjnolet @achirkin would it be possible to review the latest changes again? I've made the following changes:

Using __align__ instead of alignas in device code which clang accepts
No more weird things around __ldcv and similar functions. I simply added these intrinsics in an include which is added for CUDA compilation with clang, see the new header file under cpp/scripts
I reverted the changes around ceildiv in cluster/detail/kmeans.cuh

cjnolet

This looks good from my end. Thanks for doing this! It will be very useful once we are able to run clang-tidy across the entire codebase.

cjnolet · 2022-02-16T15:13:37Z

Also hoping to get final thoughts from @achirkin before merging.

achirkin

Looks great to me, looking forward to using clang-tidy in all cuda files!

MatthiasKohl · 2022-02-16T16:36:55Z

Thanks guys! I had to fix one final line in kmeans.cuh since my last two merges messed up the includes slightly, but CI in both RAFT and cuml should be passing now.

cjnolet · 2022-02-16T20:41:36Z

@gpucibot merge

make raft sources compilable with clang (excluding dependencies: cub,…

fd4f1b7

… cg, libcudacxx)

MatthiasKohl requested review from a team as code owners December 14, 2021 22:42

github-actions bot added the cpp label Dec 14, 2021

MatthiasKohl added bug Something isn't working CMake feature request New feature or request non-breaking Non-breaking change labels Dec 14, 2021

MatthiasKohl removed the bug Something isn't working label Dec 14, 2021

MatthiasKohl changed the title ~~[BUG] [FEA] make raft sources compilable with clang~~ [FEA] make raft sources compilable with clang Dec 14, 2021

This was referenced Dec 16, 2021

[FEA] Make cuco compilable using clang NVIDIA/cuCollections#128

Open

Feature: clang CUDA (device compilation) support NVIDIA/cccl#991

Closed

MatthiasKohl added 2 commits January 3, 2022 07:52

additional changes based on exp tests (clang13 + more compatible libc…

f25b42e

…udacxx/cuco)

Merge branch 'branch-22.02' into fea-clang-compilable

4093827

github-actions bot removed the CMake label Jan 3, 2022

cjnolet changed the base branch from branch-22.02 to branch-22.04 January 24, 2022 20:09

cjnolet requested changes Jan 24, 2022

View reviewed changes

cpp/include/raft/spectral/kmeans.hpp Outdated Show resolved Hide resolved

cpp/scripts/run-clang-compile.py Outdated Show resolved Hide resolved

cpp/scripts/run-clang-compile.py Show resolved Hide resolved

cpp/include/raft/cuda_utils.cuh Show resolved Hide resolved

MatthiasKohl added 2 commits February 11, 2022 04:03

Merge branch 'branch-22.04' into fea-clang-compilable

ecbbec3

make raft sources compatible with clang as much as possible, updated …

652fe31

…copyright

MatthiasKohl requested a review from a team as a code owner February 11, 2022 18:22

cjnolet removed the request for review from a team February 13, 2022 00:46

Merge branch 'branch-22.04' of github.com:rapidsai/raft into fea-clan…

8dbad56

…g-compilable

clang compilation: clang format fix

1ea1617

cjnolet requested changes Feb 15, 2022

View reviewed changes

achirkin reviewed Feb 15, 2022

View reviewed changes

cpp/include/raft/cuda_utils.cuh Outdated Show resolved Hide resolved

MatthiasKohl added 3 commits February 16, 2022 05:53

clang compatibility: added additional intrinsics for clang, reverted …

d4bc59f

…to use of ld/st CUDA intrinsics with cache hints

clang compatibility: using __align__ instead of alignas in device code

7b414d1

clang compatilibity: reverted ceildiv changes

dd92c9d

cjnolet approved these changes Feb 16, 2022

View reviewed changes

achirkin approved these changes Feb 16, 2022

View reviewed changes

clang compatibility: fixed includes in kmeans.cuh

608b27a

rapids-bot bot merged commit b2a88c2 into rapidsai:branch-22.04 Feb 16, 2022

vyasr mentioned this pull request Dec 21, 2022

Compile times are growing significantly rapidsai/cudf#581

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[FEA] make raft sources compilable with clang #424

[FEA] make raft sources compilable with clang #424

MatthiasKohl commented Dec 14, 2021 •

edited

Loading

MatthiasKohl commented Dec 14, 2021

cjnolet commented Dec 20, 2021

MatthiasKohl commented Dec 21, 2021

MatthiasKohl commented Jan 3, 2022

MatthiasKohl commented Jan 3, 2022

cjnolet left a comment

MatthiasKohl commented Feb 1, 2022

MatthiasKohl commented Feb 14, 2022

MatthiasKohl commented Feb 14, 2022 •

edited

Loading

cjnolet commented Feb 14, 2022

cjnolet left a comment •

edited

Loading

cjnolet Feb 14, 2022

MatthiasKohl Feb 15, 2022

MatthiasKohl commented Feb 16, 2022

cjnolet left a comment

cjnolet commented Feb 16, 2022

achirkin left a comment

MatthiasKohl commented Feb 16, 2022

cjnolet commented Feb 16, 2022

[FEA] make raft sources compilable with clang #424

[FEA] make raft sources compilable with clang #424

Conversation

MatthiasKohl commented Dec 14, 2021 • edited Loading

MatthiasKohl commented Dec 14, 2021

cjnolet commented Dec 20, 2021

MatthiasKohl commented Dec 21, 2021

MatthiasKohl commented Jan 3, 2022

MatthiasKohl commented Jan 3, 2022

cjnolet left a comment

Choose a reason for hiding this comment

MatthiasKohl commented Feb 1, 2022

MatthiasKohl commented Feb 14, 2022

MatthiasKohl commented Feb 14, 2022 • edited Loading

cjnolet commented Feb 14, 2022

cjnolet left a comment • edited Loading

Choose a reason for hiding this comment

cjnolet Feb 14, 2022

Choose a reason for hiding this comment

MatthiasKohl Feb 15, 2022

Choose a reason for hiding this comment

MatthiasKohl commented Feb 16, 2022

cjnolet left a comment

Choose a reason for hiding this comment

cjnolet commented Feb 16, 2022

achirkin left a comment

Choose a reason for hiding this comment

MatthiasKohl commented Feb 16, 2022

cjnolet commented Feb 16, 2022

MatthiasKohl commented Dec 14, 2021 •

edited

Loading

MatthiasKohl commented Feb 14, 2022 •

edited

Loading

cjnolet left a comment •

edited

Loading