Skip to content
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

[CUDA][HIP] Exclude external variables from constant promotion. #73549

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
+14 −0
Open

[CUDA][HIP] Exclude external variables from constant promotion. #73549

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions clang/lib/Sema/SemaCUDA.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -792,6 +792,7 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
(VD->isFileVarDecl() || VD->isStaticDataMember()) &&
!IsDependentVar(VD) &&
((VD->isConstexpr() || VD->getType().isConstQualified()) &&
VD->getStorageClass() != SC_Extern &&
HasAllowedCUDADeviceStaticInitializer(*this, VD,
CICK_DeviceOrConstant))) {
VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
Expand Down
13 changes: 13 additions & 0 deletions clang/test/CodeGenCUDA/device-use-host-var.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,8 @@
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s
// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s
// RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda \
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG -implicit-check-not=external_ %s

#include "Inputs/cuda.h"

Expand Down Expand Up @@ -104,3 +106,14 @@ void fun() {
(void) b<double>;
(void) var_host_only;
}

extern __global__ void external_func();
extern void* const external_dep[] = {
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This array is nomiannly host-only entity and should not be emitted on GPU at all, IMO. In fact, nvcc errors out if we attempt to access it on the GPU: https://godbolt.org/z/G15zn35Wd

Whether it's extern or not should not matter. I think.

@yxsamliu Sam, WDYT? I suspect there is/was a reason we may have allowed const access on both sides.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It seems nvcc allows non-array type const var to be used in device code but not array type const var

https://godbolt.org/z/xjvbjPK77

I don't see why we cannot use array type const var in device code if we are able to emit them on device side. There may be CUDA/HIP code already using this feature. Disabling it may cause regressions.

On the other hand, I think disallow extern const var in device code is reasonable, since we do not know how it is initialized.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sounds broken that the behavior would differ between array and non-array ?

(void*)(external_func)
};
extern void* const external_arr[] = {};

void* host_fun() {
(void) external_dep;
(void) external_arr;
}
Comment on lines +110 to +119
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There are no CHECK lines here.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This case is checked by flag "-implicit-check-not=external_" at the line 6, so we can check that external is not mentioned anywhere at the device code.