Hi pcc, rnk,
r218624 implemented target inference for implicit special members. However,
other entities can be implicit - for example intrinsics. These can not have
inference running on them, so they should be marked __host__ __device__ as
before. This is the safest and most flexible setting, since by construction
these functions don't invoke anything, and we'd like them to be invokable from
both host and device code. LLVM's intrinsics definitions (where these
intrinsics come from in the case of CUDA/NVPTX) have no notion of target, so
both host and device intrinsics can be supported this way.
http://reviews.llvm.org/D5541
Files:
lib/Sema/SemaCUDA.cpp
test/SemaCUDA/implicit-intrinsic.cu
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -48,6 +48,12 @@
if (D->hasAttr<CUDAHostAttr>())
return CFT_HostDevice;
return CFT_Device;
+ } else if (D->hasAttr<CUDAHostAttr>()) {
+ return CFT_Host;
+ } else if (D->isImplicit()) {
+ // Some implicit declarations (like intrinsic functions) are not marked.
+ // Set the most lenient target on them for maximal flexibility.
+ return CFT_HostDevice;
}
return CFT_Host;
Index: test/SemaCUDA/implicit-intrinsic.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/implicit-intrinsic.cu
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// expected-no-diagnostics
+__device__ void __threadfence_system() {
+ // This shouldn't produce an error, since __nvvm_membar_sys is inferred to
+ // be __host__ __device__ and thus callable from device code.
+ __nvvm_membar_sys();
+}
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
