From 18721a03861e3c216deae42721da3fe9ea06c9ba Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Tue, 20 Apr 2021 12:29:25 -0700 Subject: [PATCH 1/2] Add definitions for NV_TARGET_MINIMUM_* macros --- .../test/cuda/test_platform.pass.cpp | 26 ++++++++++++++ include/nv/detail/__target_macros | 36 +++++++++++++++++-- include/nv/target | 12 +++++-- 3 files changed, 69 insertions(+), 5 deletions(-) diff --git a/.upstream-tests/test/cuda/test_platform.pass.cpp b/.upstream-tests/test/cuda/test_platform.pass.cpp index cf82a0419e..404dda6bb6 100644 --- a/.upstream-tests/test/cuda/test_platform.pass.cpp +++ b/.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 diff --git a/include/nv/detail/__target_macros b/include/nv/detail/__target_macros index 11b17d2dd8..e1722f3134 100644 --- a/include/nv/detail/__target_macros +++ b/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__) diff --git a/include/nv/target b/include/nv/target index 832f1f80d0..b6998b03da 100644 --- a/include/nv/target +++ b/include/nv/target @@ -13,6 +13,14 @@ # 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 +53,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 +160,5 @@ namespace nv { } } -#endif - #include "detail/__target_macros" From d4667fa043e804c71e2dff00cadaf5420141f802 Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Wed, 28 Apr 2021 12:42:29 -0700 Subject: [PATCH 2/2] Add notice that is currently a preview --- include/nv/target | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/include/nv/target b/include/nv/target index b6998b03da..62324b41be 100644 --- a/include/nv/target +++ b/include/nv/target @@ -7,6 +7,10 @@ // //===----------------------------------------------------------------------===// +// 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__)