Add awareness for SM_87

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

@@ -36,6 +36,7 @@ __host__ __device__ void test() {
   // This test ensures that the fallthrough cases are not invoked.
   // SM_80 would imply that SM_72 is available, yet it should not be expanded by the macro
   NV_DISPATCH_TARGET(
+    NV_PROVIDES_SM_87, (static_assert(arch_val >= 870, "cuda arch expected 870");),
     NV_PROVIDES_SM_86, (static_assert(arch_val >= 860, "cuda arch expected 860");),
     NV_PROVIDES_SM_80, (static_assert(arch_val >= 800, "cuda arch expected 800");),
     NV_PROVIDES_SM_75, (static_assert(arch_val >= 750, "cuda arch expected 750");),
@@ -54,7 +55,8 @@ __host__ __device__ void test() {
 
   // This test is simpler and ensures that only the value matched is invoked, but is roughly the same as the above
   NV_DISPATCH_TARGET(
-    NV_IS_EXACTLY_SM_86, (static_assert(arch_val == 860, "cuda arch expected 800");),
+    NV_IS_EXACTLY_SM_87, (static_assert(arch_val == 870, "cuda arch expected 870");),
+    NV_IS_EXACTLY_SM_86, (static_assert(arch_val == 860, "cuda arch expected 860");),
     NV_IS_EXACTLY_SM_80, (static_assert(arch_val == 800, "cuda arch expected 800");),
     NV_IS_EXACTLY_SM_75, (static_assert(arch_val == 750, "cuda arch expected 750");),
     NV_IS_EXACTLY_SM_72, (static_assert(arch_val == 720, "cuda arch expected 720");),

include/nv/detail/__target_macros

@@ -25,6 +25,7 @@
 #  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_SELECTOR_870 nv::target::sm_87
 
 #  define _NV_TARGET_ARCH_TO_SM_350 35
 #  define _NV_TARGET_ARCH_TO_SM_370 37
@@ -39,6 +40,7 @@
 #  define _NV_TARGET_ARCH_TO_SM_750 75
 #  define _NV_TARGET_ARCH_TO_SM_800 80
 #  define _NV_TARGET_ARCH_TO_SM_860 86
+#  define _NV_TARGET_ARCH_TO_SM_870 87
 
 #if defined(_NV_COMPILER_NVCXX)
 
@@ -55,6 +57,7 @@
 #  define _NV_TARGET_VAL_SM_75 nv::target::sm_75
 #  define _NV_TARGET_VAL_SM_80 nv::target::sm_80
 #  define _NV_TARGET_VAL_SM_86 nv::target::sm_86
+#  define _NV_TARGET_VAL_SM_87 nv::target::sm_87
 
 #  define _NV_TARGET___NV_IS_HOST nv::target::is_host
 #  define _NV_TARGET___NV_IS_DEVICE nv::target::is_device
@@ -86,6 +89,7 @@
 #  define _NV_TARGET_VAL_SM_75 750
 #  define _NV_TARGET_VAL_SM_80 800
 #  define _NV_TARGET_VAL_SM_86 860
+#  define _NV_TARGET_VAL_SM_87 870
 
 #  if defined(__CUDA_ARCH__)
 #    define _NV_TARGET_VAL __CUDA_ARCH__
@@ -129,6 +133,7 @@
 #  define _NV_TARGET_VAL_SM_75 750
 #  define _NV_TARGET_VAL_SM_80 800
 #  define _NV_TARGET_VAL_SM_86 860
+#  define _NV_TARGET_VAL_SM_87 870
 
 #  define _NV_TARGET_VAL 0
 
@@ -155,6 +160,7 @@
 #define _NV_TARGET___NV_PROVIDES_SM_75 (_NV_TARGET_PROVIDES(_NV_TARGET_VAL_SM_75))
 #define _NV_TARGET___NV_PROVIDES_SM_80 (_NV_TARGET_PROVIDES(_NV_TARGET_VAL_SM_80))
 #define _NV_TARGET___NV_PROVIDES_SM_86 (_NV_TARGET_PROVIDES(_NV_TARGET_VAL_SM_86))
+#define _NV_TARGET___NV_PROVIDES_SM_87 (_NV_TARGET_PROVIDES(_NV_TARGET_VAL_SM_87))
 
 #define _NV_TARGET___NV_IS_EXACTLY_SM_35 (_NV_TARGET_IS_EXACTLY(_NV_TARGET_VAL_SM_35))
 #define _NV_TARGET___NV_IS_EXACTLY_SM_37 (_NV_TARGET_IS_EXACTLY(_NV_TARGET_VAL_SM_37))
@@ -169,6 +175,7 @@
 #define _NV_TARGET___NV_IS_EXACTLY_SM_75 (_NV_TARGET_IS_EXACTLY(_NV_TARGET_VAL_SM_75))
 #define _NV_TARGET___NV_IS_EXACTLY_SM_80 (_NV_TARGET_IS_EXACTLY(_NV_TARGET_VAL_SM_80))
 #define _NV_TARGET___NV_IS_EXACTLY_SM_86 (_NV_TARGET_IS_EXACTLY(_NV_TARGET_VAL_SM_86))
+#define _NV_TARGET___NV_IS_EXACTLY_SM_87 (_NV_TARGET_IS_EXACTLY(_NV_TARGET_VAL_SM_87))
 
 #define NV_PROVIDES_SM_35   __NV_PROVIDES_SM_35
 #define NV_PROVIDES_SM_37   __NV_PROVIDES_SM_37
@@ -183,6 +190,7 @@
 #define NV_PROVIDES_SM_75   __NV_PROVIDES_SM_75
 #define NV_PROVIDES_SM_80   __NV_PROVIDES_SM_80
 #define NV_PROVIDES_SM_86   __NV_PROVIDES_SM_86
+#define NV_PROVIDES_SM_87   __NV_PROVIDES_SM_87
 
 #define NV_IS_EXACTLY_SM_35 __NV_IS_EXACTLY_SM_35
 #define NV_IS_EXACTLY_SM_37 __NV_IS_EXACTLY_SM_37
@@ -197,6 +205,7 @@
 #define NV_IS_EXACTLY_SM_75 __NV_IS_EXACTLY_SM_75
 #define NV_IS_EXACTLY_SM_80 __NV_IS_EXACTLY_SM_80
 #define NV_IS_EXACTLY_SM_86 __NV_IS_EXACTLY_SM_86
+#define NV_IS_EXACTLY_SM_87 __NV_IS_EXACTLY_SM_87
 
 #define NV_IS_HOST         __NV_IS_HOST
 #define NV_IS_DEVICE       __NV_IS_DEVICE
@@ -290,13 +299,18 @@
 #    define _NV_TARGET_BOOL___NV_IS_EXACTLY_SM_80 0
 #  endif
 
-
 #  if (_NV_TARGET___NV_IS_EXACTLY_SM_86)
 #    define _NV_TARGET_BOOL___NV_IS_EXACTLY_SM_86 1
 #  else
 #    define _NV_TARGET_BOOL___NV_IS_EXACTLY_SM_86 0
 #  endif
 
+#  if (_NV_TARGET___NV_IS_EXACTLY_SM_87)
+#    define _NV_TARGET_BOOL___NV_IS_EXACTLY_SM_87 1
+#  else
+#    define _NV_TARGET_BOOL___NV_IS_EXACTLY_SM_87 0
+#  endif
+
 #  if (_NV_TARGET_IS_HOST)
 #    define _NV_TARGET_BOOL___NV_IS_HOST   1
 #    define _NV_TARGET_BOOL___NV_IS_DEVICE 0
@@ -388,6 +402,12 @@
 #    define _NV_TARGET_BOOL___NV_PROVIDES_SM_86 0
 #  endif
 
+#  if (_NV_TARGET___NV_PROVIDES_SM_87)
+#    define _NV_TARGET_BOOL___NV_PROVIDES_SM_87 1
+#  else
+#    define _NV_TARGET_BOOL___NV_PROVIDES_SM_87 0
+#  endif
+
 #  define _NV_ARCH_COND_CAT1(cond) _NV_TARGET_BOOL_##cond
 #  define _NV_ARCH_COND_CAT(cond) _NV_EVAL(_NV_ARCH_COND_CAT1(cond))
 
@@ -426,6 +446,8 @@
 #  define _NV_TARGET_DISPATCH_HANDLE24(q, fn, ...)  _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE22(__VA_ARGS__))
 #  define _NV_TARGET_DISPATCH_HANDLE26(q, fn, ...)  _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE24(__VA_ARGS__))
 #  define _NV_TARGET_DISPATCH_HANDLE28(q, fn, ...)  _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE26(__VA_ARGS__))
+#  define _NV_TARGET_DISPATCH_HANDLE30(q, fn, ...)  _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE28(__VA_ARGS__))
+#  define _NV_TARGET_DISPATCH_HANDLE32(q, fn, ...)  _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE30(__VA_ARGS__))
 
 #  define _NV_TARGET_DISPATCH(...) _NV_BLOCK_EXPAND(_NV_DISPATCH_N_ARY(_NV_TARGET_DISPATCH_HANDLE, __VA_ARGS__))