Page MenuHomePhabricator

[CUDA][HIP] Diagnose reference of host variable
ClosedPublic

Authored by yaxunl on Nov 11 2020, 10:02 AM.

Details

Summary

This patch diagnoses invalid references of global host variables in device,
global, or host device functions.

Diff Detail

Event Timeline

yaxunl requested review of this revision.Nov 11 2020, 10:02 AM
yaxunl created this revision.
yaxunl updated this revision to Diff 304569.Nov 11 2020, 10:11 AM

add comments about H->D not diagnosed

tra accepted this revision.Nov 11 2020, 10:12 AM

LGTM.

clang/lib/Sema/SemaExpr.cpp
357

This could use a comment why we only check D->H references.

This revision is now accepted and ready to land.Nov 11 2020, 10:12 AM
yaxunl marked an inline comment as done.Nov 11 2020, 10:14 AM
yaxunl added inline comments.
clang/lib/Sema/SemaExpr.cpp
357

Just added a comment:

Reference of device global variables in host functions is
allowed through shadow variables therefore it is not diagnosed.

This revision was automatically updated to reflect the committed changes.
yaxunl marked an inline comment as done.
Herald added a project: Restricted Project. · View Herald TranscriptDec 2 2020, 7:28 AM
Herald added a subscriber: jvesely. · View Herald Transcript
tra added a comment.Dec 8 2020, 2:26 PM

It appears that we need to add special handling for texture/surface references. Nominally they are host-side objects, but they are accessed/used from device functions as far as Sema is concerned.

E.g. https://godbolt.org/z/z1YnE3
NVCC and older clang compile it, but the recent one fails because this patch does not let the change through.
I think isCUDADeviceBuiltinTextureType has problem handling texture refs within templates.

tra added a comment.Dec 8 2020, 4:11 PM
In D91281#2441147, @tra wrote:

I think isCUDADeviceBuiltinTextureType has problem handling texture refs within templates.

Proposed fix: https://reviews.llvm.org/D92893

yaxunl added a comment.Fri, Mar 5, 9:33 AM

@tra I got some issue with this patch. There are cases that an expression using a host variable is compile-time constant, e.g.

int x;
__device__ void fun() {
  sizeof(x);
}

Do we want to allow that? Thanks.