From a6655ef2ff7f5c3e8ea46437a4c8a7bf0c5b0c48 Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Mon, 30 Aug 2021 20:37:13 -0700 Subject: [PATCH 1/3] Add support for NV_IF_TARGET in C++03 --- include/nv/detail/__preprocessor | 52 ++++++++++++++++-------- include/nv/detail/__target_macros | 66 ++++++++++++++++++++----------- include/nv/target | 2 +- 3 files changed, 79 insertions(+), 41 deletions(-) diff --git a/include/nv/detail/__preprocessor b/include/nv/detail/__preprocessor index 8f4a390ae2..27454ca933 100644 --- a/include/nv/detail/__preprocessor +++ b/include/nv/detail/__preprocessor @@ -7,6 +7,40 @@ // //===----------------------------------------------------------------------===// + +// For all compilers and dialects this header defines: +// _NV_EVAL +// _NV_IF +// _NV_CONCAT_EVAL +// For C++11 and up it defines: +// _NV_STRIP_PAREN +// _NV_DISPATCH_N_ARY +// _NV_FIRST_ARG +// _NV_REMOVE_PAREN + +#if defined(__cplusplus) && __cplusplus >= 201103L +#define _NV_EVAL2(...) __VA_ARGS__ +#define _NV_EVAL(...) _NV_EVAL2(__VA_ARGS__) +#else +#define _NV_EVAL2(x) x +#define _NV_EVAL(x) _NV_EVAL2(x) +#endif +#define _NV_CONCAT_EVAL1(l, r) _NV_EVAL(l ## r) +#define _NV_CONCAT_EVAL2(l, r) _NV_CONCAT_EVAL2(l, r) +#define _NV_CONCAT_EVAL(l, r) _NV_CONCAT_EVAL1(l, r) + +#define _NV_IF_0(t, f) f +#define _NV_IF_1(t, f) t + +#define _NV_IF_BIT(b) _NV_EVAL(_NV_IF_##b) +#define _NV_IF__EVAL(fn, t, f) _NV_EVAL(fn(t, f)) +#define _NV_IF_EVAL(cond, t, f) _NV_IF__EVAL(_NV_IF_BIT(cond), t, f) + +#define _NV_IF1(cond, t, f) _NV_IF_EVAL(cond, t, f) +#define _NV_IF(cond, t, f) _NV_IF1(_NV_EVAL(cond), _NV_EVAL(t), _NV_EVAL(f)) + +#if defined(__cplusplus) && __cplusplus >= 201103L + // The below mechanisms were derived from: https://gustedt.wordpress.com/2010/06/08/detect-empty-macro-arguments/ #define _NV_ARG32(...) _NV_EVAL(_NV_ARG32_0(__VA_ARGS__)) @@ -40,12 +74,6 @@ #define _NV_ISEMPTY0(_0, _1, _2, _3) _NV_HAS_COMMA(_NV_PASTE5(_NV_IS_EMPTY_CASE_, _0, _1, _2, _3)) #define _NV_IS_EMPTY_CASE_0001 , -#define _NV_CONCAT_EVAL1(l, r) _NV_EVAL(l ## r) -#define _NV_CONCAT_EVAL2(l, r) _NV_CONCAT_EVAL2(l, r) -#define _NV_CONCAT_EVAL(l, r) _NV_CONCAT_EVAL1(l, r) - -#define _NV_EVAL2(...) __VA_ARGS__ -#define _NV_EVAL(...) _NV_EVAL2(__VA_ARGS__) #define _NV_REMOVE_PAREN(...) _NV_REMOVE_PAREN1(__VA_ARGS__) #define _NV_REMOVE_PAREN1(...) _NV_STRIP_PAREN(_NV_IF(_NV_TEST_PAREN(__VA_ARGS__), (_NV_STRIP_PAREN(__VA_ARGS__)), (__VA_ARGS__))) @@ -71,16 +99,6 @@ #define _NV_REMOVE_FIRST_ARGS1(...) __VA_ARGS__ #define _NV_REMOVE_FIRST_ARGS(x, ...) _NV_REMOVE_FIRST_ARGS1(__VA_ARGS__) -#define _NV_IF_0(t, f) f -#define _NV_IF_1(t, ...) t - -#define _NV_IF_BIT(b) _NV_EVAL(_NV_IF_##b) -#define _NV_IF__EVAL(fn, t, f) _NV_EVAL(fn(t, f)) -#define _NV_IF_EVAL(cond, t, f) _NV_IF__EVAL(_NV_IF_BIT(cond), t, f) - -#define _NV_IF1(cond, t, f) _NV_IF_EVAL(cond, t, f) -#define _NV_IF(cond, t, f) _NV_IF1(_NV_EVAL(cond), _NV_EVAL(t), _NV_EVAL(f)) - #define _NV_NUM_ARGS(...) _NV_NUM_ARGS0(__VA_ARGS__) #define _NV_NUM_ARGS0(...) _NV_EVAL(_NV_NUM_ARGS1(__VA_ARGS__)) #define _NV_NUM_ARGS1(...) _NV_IF(_NV_ISEMPTY(__VA_ARGS__), 0, _NV_NUM_ARGS2(__VA_ARGS__)) @@ -92,3 +110,5 @@ #define _NV_DISPATCH_N_IMPL0(depth, name, ...) _NV_DISPATCH_N_IMPL1(_NV_CONCAT_EVAL(name, depth), __VA_ARGS__) #define _NV_DISPATCH_N_IMPL(name, ...) _NV_DISPATCH_N_IMPL0(_NV_NUM_ARGS(__VA_ARGS__), name, __VA_ARGS__) #define _NV_DISPATCH_N_ARY(name, ...) _NV_DISPATCH_N_IMPL(name, __VA_ARGS__) + +#endif // defined(__cplusplus) && __cplusplus >= 201103L \ No newline at end of file diff --git a/include/nv/detail/__target_macros b/include/nv/detail/__target_macros index e1722f3134..45e587669e 100644 --- a/include/nv/detail/__target_macros +++ b/include/nv/detail/__target_macros @@ -105,7 +105,7 @@ # if defined(_NV_TARGET_VAL) # define _NV_DEVICE_CHECK(q) (q) # else -# define _NV_DEVICE_CHECK(q) (false) +# define _NV_DEVICE_CHECK(q) (0) # endif # define _NV_TARGET_PROVIDES(q) _NV_DEVICE_CHECK(_NV_TARGET_VAL >= q) @@ -388,33 +388,51 @@ # define _NV_TARGET_BOOL___NV_PROVIDES_SM_86 0 # endif -# define _NV_INNER_BLOCK_EXPAND(...) __VA_ARGS__ -# define _NV_BLOCK_EXPAND(...) { _NV_REMOVE_PAREN(__VA_ARGS__) } # define _NV_ARCH_COND_CAT1(cond) _NV_TARGET_BOOL_##cond # define _NV_ARCH_COND_CAT(cond) _NV_EVAL(_NV_ARCH_COND_CAT1(cond)) -# define _NV_TARGET_IF(cond, t, ...) _NV_IF(_NV_ARCH_COND_CAT(cond), t, __VA_ARGS__) + +# if defined(__cplusplus) && __cplusplus >= 201103L + +# define _NV_BLOCK_EXPAND(...) { _NV_REMOVE_PAREN(__VA_ARGS__) } +# define _NV_TARGET_IF(cond, t, ...) _NV_IF(_NV_ARCH_COND_CAT(cond), t, __VA_ARGS__) + +# else // = 201103L + +# define _NV_TARGET_DISPATCH_HANDLE0() +# define _NV_TARGET_DISPATCH_HANDLE2(q, fn) _NV_TARGET_IF(q, fn) +# define _NV_TARGET_DISPATCH_HANDLE4(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE2(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE6(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE4(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE8(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE6(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE10(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE8(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE12(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE10(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE14(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE12(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE16(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE14(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE18(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE16(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE20(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE18(__VA_ARGS__)) +# define _NV_TARGET_DISPATCH_HANDLE22(q, fn, ...) _NV_TARGET_IF(q, fn, _NV_TARGET_DISPATCH_HANDLE20(__VA_ARGS__)) +# 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(...) _NV_BLOCK_EXPAND(_NV_DISPATCH_N_ARY(_NV_TARGET_DISPATCH_HANDLE, __VA_ARGS__)) + +# define NV_IF_TARGET(cond, t, ...) _NV_BLOCK_EXPAND(_NV_TARGET_IF(cond, t, __VA_ARGS__)) +# define NV_DISPATCH_TARGET(...) _NV_TARGET_DISPATCH(__VA_ARGS__) + +#else // = 201103L +#if defined(__cplusplus) && __cplusplus >= 201103L namespace nv { namespace target { From dfc5a5f1c04508f84ea261373025db9640eed468 Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Thu, 9 Sep 2021 17:12:32 -0700 Subject: [PATCH 2/3] Add test for c++03 mode, while not testable under NVCC, it does pass with GCC --- .../test/cuda/test_platform.pass.cpp | 2 + .../test/cuda/test_platform_cpp03.pass.cpp | 63 +++++++++++++++++++ .upstream-tests/test/force_include.h | 2 +- 3 files changed, 66 insertions(+), 1 deletion(-) create mode 100644 .upstream-tests/test/cuda/test_platform_cpp03.pass.cpp diff --git a/.upstream-tests/test/cuda/test_platform.pass.cpp b/.upstream-tests/test/cuda/test_platform.pass.cpp index 29a20cd1f9..63661cc91c 100644 --- a/.upstream-tests/test/cuda/test_platform.pass.cpp +++ b/.upstream-tests/test/cuda/test_platform.pass.cpp @@ -7,6 +7,8 @@ // //===----------------------------------------------------------------------===// +// UNSUPPORTED: c++03 + #include #if !defined(__CUDACC_RTC__) diff --git a/.upstream-tests/test/cuda/test_platform_cpp03.pass.cpp b/.upstream-tests/test/cuda/test_platform_cpp03.pass.cpp new file mode 100644 index 0000000000..cfcc3b028c --- /dev/null +++ b/.upstream-tests/test/cuda/test_platform_cpp03.pass.cpp @@ -0,0 +1,63 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: pgi, nvc++ + +#include + +#if !defined(__CUDACC_RTC__) +#include +#include +#endif + +#ifdef __CUDACC__ +# define HD_ANNO __host__ __device__ +#else +# define HD_ANNO +#endif + +template +HD_ANNO bool unused(T) {return true;} + +// Assert macro interferes with preprocessing, wrap it in a function +HD_ANNO inline void check_v(bool result) { + assert(result); +} + +HD_ANNO void test() { +# if defined(__CUDA_ARCH__) + int arch_val = __CUDA_ARCH__; +# else + int arch_val = 0; +# endif + + unused(arch_val); + + NV_IF_TARGET( + NV_IS_HOST, + check_v(arch_val == 0); + ) + + NV_IF_TARGET( + NV_IS_DEVICE, + check_v(arch_val == __CUDA_ARCH__); + ) + + NV_IF_ELSE_TARGET( + NV_IS_HOST, + check_v(arch_val == 0);, + check_v(arch_val == __CUDA_ARCH__); + ) +} + +int main(int argc, char ** argv) +{ + test(); + return 0; +} diff --git a/.upstream-tests/test/force_include.h b/.upstream-tests/test/force_include.h index b61e79e9da..a5de24a7e2 100644 --- a/.upstream-tests/test/force_include.h +++ b/.upstream-tests/test/force_include.h @@ -74,7 +74,7 @@ int main(int argc, char** argv) return ret; } - int * cuda_ret = nullptr; + int * cuda_ret = 0; CUDA_CALL(err, cudaMalloc(&cuda_ret, sizeof(int))); fake_main_kernel<<<1, cuda_thread_count>>>(cuda_ret); From 6c4f357f37394a6197451335c490536de3f520ea Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Thu, 9 Sep 2021 17:12:58 -0700 Subject: [PATCH 3/3] Formatting and naming fixes, add NV_IF_ELSE_TARGET for non-optional false statements --- include/nv/detail/__preprocessor | 12 ++++++------ include/nv/detail/__target_macros | 19 +++++++++++++------ 2 files changed, 19 insertions(+), 12 deletions(-) diff --git a/include/nv/detail/__preprocessor b/include/nv/detail/__preprocessor index 27454ca933..94acf01bae 100644 --- a/include/nv/detail/__preprocessor +++ b/include/nv/detail/__preprocessor @@ -19,14 +19,14 @@ // _NV_REMOVE_PAREN #if defined(__cplusplus) && __cplusplus >= 201103L -#define _NV_EVAL2(...) __VA_ARGS__ -#define _NV_EVAL(...) _NV_EVAL2(__VA_ARGS__) +# define _NV_EVAL1(...) __VA_ARGS__ +# define _NV_EVAL(...) _NV_EVAL1(__VA_ARGS__) #else -#define _NV_EVAL2(x) x -#define _NV_EVAL(x) _NV_EVAL2(x) -#endif +# define _NV_EVAL1(x) x +# define _NV_EVAL(x) _NV_EVAL1(x) +#endif // defined(__cplusplus) && __cplusplus >= 201103L + #define _NV_CONCAT_EVAL1(l, r) _NV_EVAL(l ## r) -#define _NV_CONCAT_EVAL2(l, r) _NV_CONCAT_EVAL2(l, r) #define _NV_CONCAT_EVAL(l, r) _NV_CONCAT_EVAL1(l, r) #define _NV_IF_0(t, f) f diff --git a/include/nv/detail/__target_macros b/include/nv/detail/__target_macros index 45e587669e..f72f55c763 100644 --- a/include/nv/detail/__target_macros +++ b/include/nv/detail/__target_macros @@ -391,15 +391,19 @@ # define _NV_ARCH_COND_CAT1(cond) _NV_TARGET_BOOL_##cond # define _NV_ARCH_COND_CAT(cond) _NV_EVAL(_NV_ARCH_COND_CAT1(cond)) +# define _NV_TARGET_EMPTY_PARAM ; + # if defined(__cplusplus) && __cplusplus >= 201103L # define _NV_BLOCK_EXPAND(...) { _NV_REMOVE_PAREN(__VA_ARGS__) } -# define _NV_TARGET_IF(cond, t, ...) _NV_IF(_NV_ARCH_COND_CAT(cond), t, __VA_ARGS__) +# define _NV_TARGET_IF(cond, t, ...) _NV_IF( _NV_ARCH_COND_CAT(cond), t, __VA_ARGS__) # else //