Merge <nv/target> updates and preview notice

Add definitions for NV_TARGET_MINIMUM_* macros
Add a notice at top of file that this is a preview feature for use in CUDA libs.

.upstream-tests/test/cuda/test_platform.pass.cpp

@@ -91,6 +91,15 @@ __host__ __device__ void test() {
       printf("Host success\r\n");,
       printf("Device success\r\n");
   )
+
+  NV_DISPATCH_TARGET(
+    NV_IS_HOST, (),
+    NV_IS_DEVICE, (
+      static_assert(NV_TARGET_MINIMUM_SM_INTEGER == (__CUDA_ARCH__ / 10), "arch mismatch");
+      static_assert(nv::target::detail::toint(NV_TARGET_MINIMUM_SM_SELECTOR) == (__CUDA_ARCH__ / 10), "arch mismatch");
+      static_assert(__CUDA_MINIMUM_ARCH__ == __CUDA_ARCH__, "arch mismatch");
+    )
+  )
 }
 
 #elif defined(TEST_NVCXX)
@@ -170,6 +179,14 @@ __host__ __device__ void test() {
       printf("Host success\r\n");,
       printf("Device success\r\n");
   )
+
+  NV_DISPATCH_TARGET(
+    NV_IS_HOST, (),
+    NV_IS_DEVICE, (
+      static_assert(NV_TARGET_MINIMUM_SM_INTEGER == (__CUDA_MINIMUM_ARCH__ / 10), "arch mismatch");
+      static_assert(nv::target::detail::toint(NV_TARGET_MINIMUM_SM_SELECTOR) == (__CUDA_MINIMUM_ARCH__ / 10), "arch mismatch");
+    )
+  )
 }
 
 #elif defined(TEST_HOST)
@@ -240,6 +257,15 @@ void test() {
       printf("Host success\r\n");,
       printf("Device success\r\n");
   )
+
+  NV_DISPATCH_TARGET(
+    NV_IS_HOST, (),
+    NV_IS_DEVICE, (
+      static_assert(NV_TARGET_MINIMUM_SM_INTEGER == (__CUDA_ARCH__ / 10), "arch mismatch");
+      static_assert(nv::target::detail::toint(NV_TARGET_MINIMUM_SM_SELECTOR) == (__CUDA_ARCH__ / 10), "arch mismatch");
+      static_assert(__CUDA_MINIMUM_ARCH__ == __CUDA_ARCH__, "arch mismatch");
+    )
+  )
 }
 
 #endif

include/nv/detail/__target_macros

@@ -12,6 +12,34 @@
 
 #include "__preprocessor"
 
+#  define _NV_TARGET_ARCH_TO_SELECTOR_350 nv::target::sm_35
+#  define _NV_TARGET_ARCH_TO_SELECTOR_370 nv::target::sm_37
+#  define _NV_TARGET_ARCH_TO_SELECTOR_500 nv::target::sm_50
+#  define _NV_TARGET_ARCH_TO_SELECTOR_520 nv::target::sm_52
+#  define _NV_TARGET_ARCH_TO_SELECTOR_530 nv::target::sm_53
+#  define _NV_TARGET_ARCH_TO_SELECTOR_600 nv::target::sm_60
+#  define _NV_TARGET_ARCH_TO_SELECTOR_610 nv::target::sm_61
+#  define _NV_TARGET_ARCH_TO_SELECTOR_620 nv::target::sm_62
+#  define _NV_TARGET_ARCH_TO_SELECTOR_700 nv::target::sm_70
+#  define _NV_TARGET_ARCH_TO_SELECTOR_720 nv::target::sm_72
+#  define _NV_TARGET_ARCH_TO_SELECTOR_750 nv::target::sm_75
+#  define _NV_TARGET_ARCH_TO_SELECTOR_800 nv::target::sm_80
+#  define _NV_TARGET_ARCH_TO_SELECTOR_860 nv::target::sm_86
+
+#  define _NV_TARGET_ARCH_TO_SM_350 35
+#  define _NV_TARGET_ARCH_TO_SM_370 37
+#  define _NV_TARGET_ARCH_TO_SM_500 50
+#  define _NV_TARGET_ARCH_TO_SM_520 52
+#  define _NV_TARGET_ARCH_TO_SM_530 53
+#  define _NV_TARGET_ARCH_TO_SM_600 60
+#  define _NV_TARGET_ARCH_TO_SM_610 61
+#  define _NV_TARGET_ARCH_TO_SM_620 62
+#  define _NV_TARGET_ARCH_TO_SM_700 70
+#  define _NV_TARGET_ARCH_TO_SM_720 72
+#  define _NV_TARGET_ARCH_TO_SM_750 75
+#  define _NV_TARGET_ARCH_TO_SM_800 80
+#  define _NV_TARGET_ARCH_TO_SM_860 86
+
 #if defined(_NV_COMPILER_NVCXX)
 
 #  define _NV_TARGET_VAL_SM_35 nv::target::sm_35
@@ -35,7 +63,9 @@
 #  define _NV_TARGET___NV_NO_TARGET (nv::target::no_target)
 
 #  if defined(NV_TARGET_SM_INTEGER_LIST)
-#    define _NV_TARGET_MINIMUM_VAL (_NV_FIRST_ARG(NV_TARGET_SM_INTEGER_LIST) * 10)
+#    define NV_TARGET_MINIMUM_SM_SELECTOR _NV_FIRST_ARG(NV_TARGET_SM_SELECTOR_LIST)
+#    define NV_TARGET_MINIMUM_SM_INTEGER _NV_FIRST_ARG(NV_TARGET_SM_INTEGER_LIST)
+#    define __CUDA_MINIMUM_ARCH__ _NV_CONCAT_EVAL(_NV_FIRST_ARG(NV_TARGET_SM_INTEGER_LIST), 0)
 #  endif
 
 #  define _NV_TARGET_PROVIDES(q)   nv::target::provides(q)
@@ -59,7 +89,9 @@
 
 #  if defined(__CUDA_ARCH__)
 #    define _NV_TARGET_VAL __CUDA_ARCH__
-#    define _NV_TARGET_MINIMUM_VAL __CUDA_ARCH__
+#    define NV_TARGET_MINIMUM_SM_SELECTOR _NV_CONCAT_EVAL(_NV_TARGET_ARCH_TO_SELECTOR_, __CUDA_ARCH__)
+#    define NV_TARGET_MINIMUM_SM_INTEGER _NV_CONCAT_EVAL(_NV_TARGET_ARCH_TO_SM_, __CUDA_ARCH__)
+#    define __CUDA_MINIMUM_ARCH__ __CUDA_ARCH__
 #  endif
 
 #  if defined(__CUDA_ARCH__)

include/nv/target

@@ -7,12 +7,24 @@
 //
 //===----------------------------------------------------------------------===//
 
+// This header contains a preview of a portability system that enables
+// CUDA C++ development with NVC++, NVCC, and supported host compilers.
+// These interfaces are not guaranteed to be stable.
+
 #pragma once
 
 #if defined(__NVCC__) || defined(__CUDACC_RTC__)
 #  define _NV_COMPILER_NVCC
 #elif defined(__NVCOMPILER)
 #  define _NV_COMPILER_NVCXX
+#else
+#endif
+
+#if defined(_NV_COMPILER_NVCXX)
+#  define _NV_BITSET_ATTRIBUTE [[nv::__target_bitset]]
+#else
+#  define _NV_BITSET_ATTRIBUTE
+#endif
 
 namespace nv {
   namespace target {
@@ -45,7 +57,7 @@ namespace nv {
           sm_80_bit | sm_86_bit;
 
       // Store a set of targets as a set of bits
-      struct [[nv::__target_bitset]] target_description {
+      struct _NV_BITSET_ATTRIBUTE target_description {
         base_int_t targets;
         constexpr target_description(base_int_t a) : targets(a) { }
       };
@@ -152,7 +164,5 @@ namespace nv {
   }
 }
 
-#endif
-
 #include "detail/__target_macros"