Skip to content

need to add -Wno-nvcc-compt #1330

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Closed
yxsamliu opened this issue Jun 11, 2024 · 7 comments
Closed

need to add -Wno-nvcc-compt #1330

yxsamliu opened this issue Jun 11, 2024 · 7 comments
Labels
urgency_blocker blocking feature deliverables

Comments

@yxsamliu
Copy link

yxsamliu commented Jun 11, 2024 •
edited
Loading

A recent clang change (llvm/llvm-project#77359) caused build failure in composable_kernel:

[2024-06-11T01:46:46.213Z] /jenkins/workspace/compiler-psdb-amd-staging/Libs/composable_kernel/include/ck/utility/f8_utils.hpp:20:23: error: target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:new declaration is __device__ function, old declaration is __host__ function [-Werror,-Wnvcc-compat]
[2024-06-11T01:46:46.213Z] 20 | __device__ inline int clz(uint32_t x) { return __clz(x); }
[2024-06-11T01:46:46.213Z] | ^
[2024-06-11T01:46:46.213Z] /jenkins/workspace/compiler-psdb-amd-staging/Libs/composable_kernel/include/ck/utility/f8_utils.hpp:19:21: note: previous declaration is here
[2024-06-11T01:46:46.213Z] 19 | __host__ inline int clz(uint32_t x) { return __builtin_clz(x); }
[2024-06-11T01:46:46.213Z] |

This warning is intended for HIP code that is compatible with nvcc (not overloading device/host functions with the same signature) and is off by default. composable_kernel uses overloaded device/host functions with the same signature, therefore has to turn this warning off. Since composable_kernel uses -Weverything, it has to explicitly add -Wno-nvcc-compat to disable this warning.
@yxsamliu
Copy link
Author

@junliume This is blocking a compiler feature requested by customers. Can some one take a look? Thanks.

@bartekxk bartekxk mentioned this issue Jun 14, 2024
Merged
@bartekxk
Copy link
Contributor

Hi @yxsamliu , could you verify with this PR #1342 ?

@aosewski
Copy link
Collaborator

@yxsamliu any update? Does it work now?

@yxsamliu
Copy link
Author

I will try it and get back to you. thanks

@yxsamliu
Copy link
Author

I can verify the PR fixes the issue.

@bartekxk
Copy link
Contributor

Hi @yxsamliu I built llvm-project and Im able to reproduce this bug. I will try to resolve them.

@junliume junliume added the urgency_blocker blocking feature deliverables label Jun 18, 2024
@aosewski
Copy link
Collaborator

Fix merged into develop:#1342

searlmc1 pushed a commit to ROCm/llvm-project that referenced this issue Aug 22, 2024
@yxsamliu
[CUDA][HIP] warn incompatible redeclare (llvm#77359)
bf3cf4d 
This PR was already merged in trunk:

llvm#77359

However, it was reverted in amd-staging due
to ck issue ROCm/composable_kernel#1330

Now try relanding it in amd-staging after ck issue fixed
by

ROCm/composable_kernel#1342

Fixes: SWDEV-431838

Original commit message:

nvcc warns about the following code:

`void f();
__device__ void f() {}`

but clang does not since clang allows device function to overload host
function.

Users want clang to emit similar warning to help code to be compatible
with nvcc.

Since this may cause regression with existing code, the warning is off
by default and can be enabled by -Wnvcc-compat.

It won't cause warning in system headers, even with -Wnvcc-compat.

Change-Id: Ia370700eb3eb1b1928d04ec59e2ec63506f85545
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
urgency_blocker blocking feature deliverables
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
4 participants
@yxsamliu @junliume @aosewski @bartekxk