Revert shenanigans for launching kernels with clang.

The code being reverted compiles, but doesn't actually work -- in fact,
we end up generating no code on the device at all!

The problem is that triple_chevron_launcher_base is a templated class,
and so we only instantiate functions that are actually called.  When
compiling for device, we would never instantiate the __host__ function,
so we never observed the kernel being referenced, so we never generated
any code for it.

clang was actually wrong to reject this code.  If the launch functions
were __device__, clang would have been right.  But being __host__
__device__, clang should wait to emit an error until it sees the HD
function being *used* on the device side.  If the function is never used
on the device side, it's fine to do any host-only things from the HD
function.

As of r283963 [1], clang now does the right thing.

[1] llvm-mirror/clang@f3d02c1

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

@@ -71,27 +71,11 @@ struct triple_chevron_launcher_base<block_size,Function,true>
 {
   typedef void (*global_function_pointer_t)(Function);
 
-#if defined(__CUDA__) && defined(__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
 };
 
 
@@ -111,27 +95,11 @@ struct triple_chevron_launcher_base<block_size,Function,false>
 {
   typedef void (*global_function_pointer_t)(const Function*);
 
-#if defined (__CUDA__) && defined(__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
 };
 
 
@@ -241,3 +209,4 @@ class triple_chevron_launcher<block_size_,Function,false> : protected triple_che
 } // end detail
 } // end bul
 BULK_NAMESPACE_SUFFIX
+

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

@@ -75,27 +75,11 @@ template<typename Closure,
 {
   typedef void (*launch_function_t)(Closure); 
 
-#if defined(__CUDA__) && defined(__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__
@@ -132,27 +116,11 @@ template<typename Closure>
 {
   typedef void (*launch_function_t)(const Closure *); 
 
-#if defined(__CUDA__) && defined(__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__