Merge pull request #204 from NVIDIA/bugfix/nv_target_c98

Create macro fallback support for `<nv/target>` in C++03 and C

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

@@ -7,6 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+// UNSUPPORTED: c++03
+
 #include <nv/target>
 
 #if !defined(__CUDACC_RTC__)

.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 <nv/target>
+
+#if !defined(__CUDACC_RTC__)
+#include <assert.h>
+#include <stdio.h>
+#endif
+
+#ifdef __CUDACC__
+# define HD_ANNO __host__ __device__
+#else
+# define HD_ANNO
+#endif
+
+template <typename T>
+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;
+}

.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);

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_EVAL1(...) __VA_ARGS__
+#  define _NV_EVAL(...) _NV_EVAL1(__VA_ARGS__)
+#else
+#  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_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

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,58 @@
 #    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__)
+
+#    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__)
+
+#  else // <C++11 fallback
+
+#    define _NV_BLOCK_EXPAND(x) { x }
+
+#    define _NV_TARGET_IF(cond, t)         _NV_IF(_NV_ARCH_COND_CAT(cond), t, _NV_TARGET_EMPTY_PARAM)
+#    define _NV_TARGET_IF_ELSE(cond, t, f) _NV_IF(_NV_ARCH_COND_CAT(cond), t, f)
+
+#  endif
 
 #endif // _NV_COMPILER_NVCC
 
-#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__)
+#if defined(__cplusplus) && __cplusplus >= 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__))
+
+// NV_IF_TARGET supports a false statement provided as a variadic macro
+#  define NV_IF_TARGET(cond, t, ...)    _NV_BLOCK_EXPAND(_NV_TARGET_IF(cond, t, __VA_ARGS__))
+#  define NV_IF_ELSE_TARGET(cond, t, f) _NV_BLOCK_EXPAND(_NV_TARGET_IF(cond, t, f))
+#  define NV_DISPATCH_TARGET(...)       _NV_TARGET_DISPATCH(__VA_ARGS__)
+
+#else // <C++11 fallback
+
+// NV_IF_TARGET does not support a fallback false statement in C++03 or C dialects
+#  define NV_IF_TARGET(cond, t)         _NV_BLOCK_EXPAND(_NV_TARGET_IF(cond, t))
+#  define NV_IF_ELSE_TARGET(cond, t, f) _NV_BLOCK_EXPAND(_NV_TARGET_IF_ELSE(cond, t, f))
+
+#endif
 
 #endif // _NV__TARGET_MACROS

include/nv/target

@@ -31,7 +31,7 @@
 #  define _NV_BITSET_ATTRIBUTE
 #endif
 
-#if __cplusplus >= 201103L
+#if defined(__cplusplus) && __cplusplus >= 201103L
 
 namespace nv {
   namespace target {