Page MenuHomePhabricator

[CUDA] Do not diagnose host/device variable access in dependent types.

Authored by tra on Dec 8 2020, 3:17 PM.



isCUDADeviceBuiltinSurfaceType()/isCUDADeviceBuiltinTextureType() do not
work on dependent types as they rely on specific type attributes and that
results in mis-diagnosing access when the check happens on class templates used as texture references.

Diff Detail

Event Timeline

tra created this revision.Dec 8 2020, 3:17 PM
tra requested review of this revision.Dec 8 2020, 3:17 PM
Herald added a project: Restricted Project. · View Herald TranscriptDec 8 2020, 3:17 PM
yaxunl added a comment.Dec 8 2020, 5:14 PM

LGTM. Can we have a test?

hliao added a comment.Dec 8 2020, 11:40 PM

LGTM if there's a regression test available.

tra updated this revision to Diff 311065.Dec 10 2020, 4:41 PM
tra edited the summary of this revision. (Show Details)

Found another corner case (reference within a template with the surface/texture
attibute.) and figured out a better fix.

Added a test case.

tra added a subscriber: rsmith.Dec 10 2020, 4:49 PM

My first variant of the patch only helped with some cases when the surface/texture attribute type was used.
Trying to reduce real-world failure resulted in an example that I've added as the test case which was still failing with this patch applied.

Thanks to @rsmith, who suggested applying MeaningfulToClassTemplateDefinition, the attributes should work for templates now, too.

yaxunl added inline comments.Dec 10 2020, 8:11 PM

use __device__ will make it more readable. same as below.

hliao accepted this revision.Dec 10 2020, 9:34 PM

LGTM if you revise the test based on Sam's suggestion on the test case.

This revision is now accepted and ready to land.Dec 10 2020, 9:34 PM
tra updated this revision to Diff 311657.Dec 14 2020, 11:09 AM

Use device in the test case.

tra marked an inline comment as done.Dec 14 2020, 11:09 AM
This revision was landed with ongoing or failed builds.Dec 14 2020, 11:56 AM
This revision was automatically updated to reflect the committed changes.