This is an archive of the discontinued LLVM Phabricator instance.

[CUDA][HIP] Allow non-ODR use of host var in device
ClosedPublic

Authored by yaxunl on Mar 8 2021, 9:36 AM.

Diff Detail

Event Timeline

yaxunl requested review of this revision.Mar 8 2021, 9:36 AM
yaxunl created this revision.
tra added inline comments.Mar 8 2021, 11:30 AM
clang/test/SemaCUDA/device-use-host-var.cu
55

I do not think it should be allowed. We end up instantiating the variable on device, even though the variable should be host-only.

Right now we allow it, but end up with an .extern .const which will make ptxas fail:
https://godbolt.org/z/sx9845

If we do allow it, we'll need to make sure that we only use the value, but do not allow instantiating the variable.

yaxunl marked an inline comment as done.Mar 10 2021, 5:44 AM
yaxunl added inline comments.
clang/test/SemaCUDA/device-use-host-var.cu
55

will only not allow this since it results in ODR-use of the host var.

yaxunl updated this revision to Diff 329632.Mar 10 2021, 5:47 AM
yaxunl marked an inline comment as done.

Follow C++ about ODR-use of variables.

yaxunl updated this revision to Diff 329686.Mar 10 2021, 9:04 AM

minor bug fix

tra added a reviewer: rsmith.Mar 10 2021, 10:11 AM
tra added a subscriber: rsmith.

LGTM

I've added @rsmith to double check that we're handling it correctly.

tra added a comment.Mar 29 2021, 10:30 AM

I think we also may want to check that we allow sizeof(host_var) in the GPU code.

In D98193#2656326, @tra wrote:

I think we also may want to check that we allow sizeof(host_var) in the GPU code.

We have tests for that at line 94 of test/SemaCUDA/device-use-host-var.cu

rsmith added inline comments.Mar 29 2021, 12:31 PM
clang/lib/Sema/SemaExpr.cpp
17050

I suspect you want hasGlobalStorage rather than isFileVarDecl here (that is, preserve the condition from the deleted code), in order to disallow use of host-side local static variables from device-side functions:

__host__ void f() {
  static int n;
  struct X {
    __device__ void g() { ++n; }
  };
  // ...
}

In principle this seems reasonable to me.

yaxunl added inline comments.Apr 6 2021, 12:04 PM
clang/lib/Sema/SemaExpr.cpp
17050

For function scope static variable, if the smallest enclosing function is device or device host function, the static variable without __device__ or __constant__ attribute is allowed and the variable is emitted at device side (https://godbolt.org/z/PY5d3WGas). In that case, a device function is allowed to access that static variable even if it does not have __device__ or __constant__ attribute.

I will make changes to handle the function scope static variable.

yaxunl updated this revision to Diff 335702.Apr 6 2021, 7:51 PM

revised by Richard's comments. Check function-scope static var.

yaxunl updated this revision to Diff 335904.Apr 7 2021, 12:31 PM

fix test failure on windows. need to specify triple since it affects name mangling.

tra accepted this revision.Apr 19 2021, 9:26 AM

Small test nit. LGTM otherwise.

clang/test/CodeGenCUDA/device-use-host-var.cu
23

Nit: You may want to add a CHECK-LABEL: <function name>. It does not make much of a difference in this file, but is useful in general to explicitly restrict the search scope.

This revision is now accepted and ready to land.Apr 19 2021, 9:26 AM
yaxunl marked an inline comment as done.Apr 19 2021, 9:29 AM
yaxunl added inline comments.
clang/test/CodeGenCUDA/device-use-host-var.cu
23

will do

yaxunl updated this revision to Diff 338552.Apr 19 2021, 10:40 AM
yaxunl marked an inline comment as done.

Rebase and fix HIP header bug exposed by this patch.

tra accepted this revision.Apr 19 2021, 10:52 AM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptApr 19 2021, 11:45 AM