Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

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

Merged
jaredhoberock merged 1 commit into from
Sep 14, 2016
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 12 additions & 0 deletions thrust/system/cuda/detail/bulk/detail/cuda_launcher/triple_chevron_launcher.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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__)
Copy link
Contributor

@3gx 3gx Sep 14, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The correct condition should be

#if !defined(__CUDA_ARCH__) || (defined(__CUDACC_RDC__) && __CUDA_ARCH >= 350)

return launch_by_value<block_size,Function>;
#else
return NULL;
#endif
}
};

Expand All @@ -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
}
};

Expand Down
13 changes: 13 additions & 0 deletions thrust/system/cuda/detail/detail/launch_closure.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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>
Expand Down Expand Up @@ -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>
Expand Down