Add support for sm_90a in <nv/target> API (#1411)

* Add SM90a support to <nv/target>

Can be used as follows:

__global__ void kernel(int *out) {
    NV_DISPATCH_TARGET(
        NV_HAS_FEATURE_SM_90a, (
        *out = 901;
    ), NV_PROVIDES_SM_90, (
        *out = 900;
    ), NV_IS_DEVICE, (
        *out = 0;
    ));
}

* Fix nvc++ support

libcudacxx/include/nv/detail/__target_macros

@@ -231,6 +231,10 @@
 #define NV_IS_EXACTLY_SM_89 __NV_IS_EXACTLY_SM_89
 #define NV_IS_EXACTLY_SM_90 __NV_IS_EXACTLY_SM_90
 
+// Disable SM_90a support on non-supporting compilers.
+// Will re-enable for nvcc below.
+#define NV_HAS_FEATURE_SM_90a NV_NO_TARGET
+
 #define NV_IS_HOST         __NV_IS_HOST
 #define NV_IS_DEVICE       __NV_IS_DEVICE
 
@@ -347,6 +351,15 @@
 #    define _NV_TARGET_BOOL___NV_IS_EXACTLY_SM_90 0
 #  endif
 
+// Re-enable sm_90a support in nvcc.
+#  undef NV_HAS_FEATURE_SM_90a
+#  define NV_HAS_FEATURE_SM_90a   __NV_HAS_FEATURE_SM_90a
+#  if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900) && defined(__CUDA_ARCH_FEAT_SM90_ALL))
+#    define _NV_TARGET_BOOL___NV_HAS_FEATURE_SM_90a 1
+#  else
+#    define _NV_TARGET_BOOL___NV_HAS_FEATURE_SM_90a 0
+#  endif
+
 #  if (_NV_TARGET_IS_HOST)
 #    define _NV_TARGET_BOOL___NV_IS_HOST   1
 #    define _NV_TARGET_BOOL___NV_IS_DEVICE 0