Merge pull request #842 from andrewcorrigan/issue_385_or_how_about_ov…

…erloading_in_clang use overloading with clang to avoid taking a reference to a __global_…
NVIDIA · Oct 11, 2016 · 7a8ea01 · 7a8ea01
2 parents 2f9ae57 + 32066a2
commit 7a8ea01
Show file tree

Hide file tree

Showing 3 changed files with 67 additions and 3 deletions.
diff --git a/thrust/detail/config/compiler.h b/thrust/detail/config/compiler.h
@@ -71,10 +71,10 @@
 #define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_GCC
 #elif THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_CLANG
 // CUDA-capable clang should behave similar to NVCC.
-#if defined(__CUDA__)
-#define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_NVCC
-#else
+#if defined(__CUDA__) && defined(__clang__)
 #define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_CLANG
+#else
+#define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_NVCC
 #endif
 #else
 #define THRUST_DEVICE_COMPILER THRUST_DEVICE_COMPILER_UNKNOWN

diff --git a/thrust/system/cuda/detail/bulk/detail/cuda_launcher/triple_chevron_launcher.hpp b/thrust/system/cuda/detail/bulk/detail/cuda_launcher/triple_chevron_launcher.hpp
@@ -71,11 +71,27 @@ struct triple_chevron_launcher_base<block_size,Function,true>
 {
   typedef void (*global_function_pointer_t)(Function);
 
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_CLANG
+  __host__
+  static global_function_pointer_t global_function_pointer()
+  {
+    return launch_by_value<block_size,Function>;
+  }
+
+  __device__
+  static global_function_pointer_t global_function_pointer()
+  {
+    bulk::detail::terminate(); //  clang doesn't support dynamic parallelism
+
+    return NULL;
+  }
+#else
   __host__ __device__
   static global_function_pointer_t global_function_pointer()
   {
     return launch_by_value<block_size,Function>;
   }
+#endif
 };
 
 
@@ -95,11 +111,27 @@ struct triple_chevron_launcher_base<block_size,Function,false>
 {
   typedef void (*global_function_pointer_t)(const Function*);
 
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_CLANG
+  __host__
+  static global_function_pointer_t global_function_pointer()
+  {
+    return launch_by_pointer<block_size,Function>;
+  }
+
+  __device__
+  static global_function_pointer_t global_function_pointer()
+  {
+    bulk::detail::terminate(); //  clang doesn't support dynamic parallelism
+
+    return NULL;
+  }
+#else
   __host__ __device__
   static global_function_pointer_t global_function_pointer()
   {
     return launch_by_pointer<block_size,Function>;
   }
+#endif
 };
 
 

diff --git a/thrust/system/cuda/detail/detail/launch_closure.inl b/thrust/system/cuda/detail/detail/launch_closure.inl
@@ -75,11 +75,27 @@ template<typename Closure,
 {
   typedef void (*launch_function_t)(Closure); 
 
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_CLANG
+  __host__
+  static launch_function_t get_launch_function()
+  {
+    return launch_closure_by_value<Closure>;
+  }
+
+  __device__
+  static launch_function_t get_launch_function()
+  {
+    assert(0); // clang doesn't support dynamic parallelism
+
+    return NULL;
+  }
+#else
   __host__ __device__
   static launch_function_t get_launch_function()
   {
     return launch_closure_by_value<Closure>;
   }
+#endif
 
   template<typename DerivedPolicy, typename Size1, typename Size2, typename Size3>
   __host__ __device__
@@ -116,11 +132,27 @@ template<typename Closure>
 {
   typedef void (*launch_function_t)(const Closure *); 
 
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_CLANG
+  __host__
+  static launch_function_t get_launch_function(void)
+  {
+    return launch_closure_by_pointer<Closure>;
+  }
+
+  __device__
+  static launch_function_t get_launch_function(void)
+  {
+    assert(0); // clang doesn't support dynamic parallelism
+
+    return NULL;
+  }
+#else
   __host__ __device__
   static launch_function_t get_launch_function(void)
   {
     return launch_closure_by_pointer<Closure>;
   }
+#endif
 
   template<typename DerivedPolicy, typename Size1, typename Size2, typename Size3>
   __host__ __device__