Merge pull request #831 from jlebar/launch-host-only

Mark functions that take the address of a __global__ function as host…

Author: jaredhoberock

thrust/system/cuda/detail/bulk/detail/cuda_launcher/triple_chevron_launcher.hpp

@@ -74,7 +74,13 @@ struct triple_chevron_launcher_base<block_size,Function,true>
   __host__ __device__
   static global_function_pointer_t global_function_pointer()
   {
+    // Don't try to take the address of launch_by_value from the device side if
+    // we don't support launching kernels from __device__ functions.
+#if !defined(__CUDA_ARCH__) || defined(__CUDACC_RDC__)
     return launch_by_value<block_size,Function>;
+#else
+    return NULL;
+#endif
   }
 };
 
@@ -98,7 +104,13 @@ struct triple_chevron_launcher_base<block_size,Function,false>
   __host__ __device__
   static global_function_pointer_t global_function_pointer()
   {
+    // Don't try to take the address of launch_by_pointer from the device side
+    // if we don't support launching kernels from __device__ functions.
+#if !defined(__CUDA_ARCH__) || defined(__CUDACC_RDC__)
     return launch_by_pointer<block_size,Function>;
+#else
+    return NULL;
+#endif
   }
 };
 

thrust/system/cuda/detail/detail/launch_closure.inl

@@ -78,7 +78,13 @@ template<typename Closure,
   __host__ __device__
   static launch_function_t get_launch_function()
   {
+    // Don't try to take the address of launch_closure_by_value from the device
+    // side if we don't support launching kernels from __device__ functions.
+#if !defined(__CUDA_ARCH__) || defined(__CUDACC_RDC__)
     return launch_closure_by_value<Closure>;
+#else
+    return NULL;
+#endif
   }
 
   template<typename DerivedPolicy, typename Size1, typename Size2, typename Size3>
@@ -119,7 +125,14 @@ template<typename Closure>
   __host__ __device__
   static launch_function_t get_launch_function(void)
   {
+    // Don't try to take the address of launch_closure_by_pointer from the
+    // device side if we don't support launching kernels from __device__
+    // functions.
+#if !defined(__CUDA_ARCH__) || defined(__CUDACC_RDC__)
     return launch_closure_by_pointer<Closure>;
+#else
+    return NULL;
+#endif
   }
 
   template<typename DerivedPolicy, typename Size1, typename Size2, typename Size3>