2424# define _NV_COMPILER_CLANG_CUDA
2525#endif
2626
27- #if defined(__CUDACC_RTC__)
28- # define _NV_FUNCTION_ANNOTATION __device__
29- #else
30- # define _NV_FUNCTION_ANNOTATION
31- #endif
32-
3327#if defined(_NV_COMPILER_NVCXX)
3428# define _NV_BITSET_ATTRIBUTE [[nv::__target_bitset]]
3529#else
4236# define _NV_TARGET_CPP11
4337#endif
4438
45- #if defined(_NV_TARGET_CPP11)
39+ // Hide `if target` support from NVCC/NVRTC
40+ #if defined(_NV_TARGET_CPP11) && !defined(_NV_COMPILER_NVCC)
4641
4742namespace nv {
4843 namespace target {
@@ -81,7 +76,7 @@ namespace nv {
8176 // Store a set of targets as a set of bits
8277 struct _NV_BITSET_ATTRIBUTE target_description {
8378 base_int_t targets;
84- _NV_FUNCTION_ANNOTATION
79+
8580 constexpr target_description(base_int_t a) : targets(a) { }
8681 };
8782
@@ -94,11 +89,11 @@ namespace nv {
9489 sm_80 = 80, sm_86 = 86, sm_87 = 87,
9590 sm_89 = 89, sm_90 = 90,
9691 };
97- _NV_FUNCTION_ANNOTATION
92+
9893 constexpr base_int_t toint(sm_selector a) {
9994 return static_cast<base_int_t>(a);
10095 }
101- _NV_FUNCTION_ANNOTATION
96+
10297 constexpr base_int_t bitexact(sm_selector a) {
10398 return toint(a) == 35 ? sm_35_bit :
10499 toint(a) == 37 ? sm_37_bit :
@@ -117,7 +112,7 @@ namespace nv {
117112 toint(a) == 89 ? sm_89_bit :
118113 toint(a) == 90 ? sm_90_bit : 0;
119114 }
120- _NV_FUNCTION_ANNOTATION
115+
121116 constexpr base_int_t bitrounddown(sm_selector a) {
122117 return toint(a) >= 90 ? sm_90_bit :
123118 toint(a) >= 89 ? sm_89_bit :
@@ -139,31 +134,26 @@ namespace nv {
139134
140135 // Public API for NVIDIA GPUs
141136
142- _NV_FUNCTION_ANNOTATION
143137 constexpr target_description is_exactly(sm_selector a) {
144138 return target_description(bitexact(a));
145139 }
146140
147- _NV_FUNCTION_ANNOTATION
148141 constexpr target_description provides(sm_selector a) {
149142 return target_description(~(bitrounddown(a) - 1) & all_devices);
150143 }
151144
152145 // Boolean operations on target sets
153146
154- _NV_FUNCTION_ANNOTATION
155147 constexpr target_description operator&&(target_description a,
156148 target_description b) {
157149 return target_description(a.targets & b.targets);
158150 }
159151
160- _NV_FUNCTION_ANNOTATION
161152 constexpr target_description operator||(target_description a,
162153 target_description b) {
163154 return target_description(a.targets | b.targets);
164155 }
165156
166- _NV_FUNCTION_ANNOTATION
167157 constexpr target_description operator!(target_description a) {
168158 return target_description(~a.targets & (all_devices | all_hosts));
169159 }
@@ -205,7 +195,7 @@ namespace nv {
205195 }
206196}
207197
208- #endif // C++11
198+ #endif // C++11 && !defined(_NV_COMPILER_NVCC)
209199
210200#include "detail/__target_macros"
211201