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

Initial port to NV_IF_TARGET. #1404

Closed
alliepiper wants to merge 8 commits into from
Closed

Initial port to NV_IF_TARGET. #1404

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
824452e
Bump CUB for NV_IF_TARGET refactor.
alliepiper Mar 26, 2021
52046e8
Remove checks for obsolete architectures.
alliepiper Mar 26, 2021
5a32128
Refactor to use NV_IF_TARGET.
alliepiper Mar 26, 2021
fdc96c5
Port AgentLauncher to use cub::detail::ptx_dispatch.
alliepiper May 6, 2021
eac1b72
Use CUB's new CDP macros.
alliepiper May 25, 2021
2e102df
Remove unreachable code.
alliepiper May 27, 2021
aecdc61
Initialize members in `cuda_optional` detail class.
alliepiper May 28, 2021
f0d37ca
Fix some new and exciting exec_space `[subobject]` warnings.
alliepiper Jul 29, 2021
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
Prev Previous commit
Next Next commit
Remove checks for obsolete architectures.
  • Loading branch information
@alliepiper
alliepiper committed Jan 24, 2022
commit 52046e84768c7bd3a0c7080c0e3c7e4083b3094b
4 changes: 0 additions & 4 deletions testing/cuda/pair_sort.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,12 +8,8 @@ template<typename ExecutionPolicy, typename Iterator1, typename Iterator2>
__global__
void stable_sort_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 is_supported)
{
#if (__CUDA_ARCH__ >= 200)
*is_supported = true;
thrust::stable_sort(exec, first, last);
#else
*is_supported = false;
#endif
}


Expand Down
4 changes: 0 additions & 4 deletions testing/cuda/pair_sort_by_key.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,12 +10,8 @@ template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typen
__global__
void stable_sort_by_key_kernel(ExecutionPolicy exec, Iterator1 keys_first, Iterator1 keys_last, Iterator2 values_first, Iterator3 is_supported)
{
#if (__CUDA_ARCH__ >= 200)
*is_supported = true;
thrust::stable_sort_by_key(exec, keys_first, keys_last, values_first);
#else
*is_supported = false;
#endif
}


Expand Down
8 changes: 0 additions & 8 deletions testing/cuda/partition.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -290,12 +290,8 @@ template<typename ExecutionPolicy, typename Iterator1, typename Predicate, typen
__global__
void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Predicate pred, Iterator2 result, Iterator3 is_supported)
{
#if (__CUDA_ARCH__ >= 200)
*is_supported = true;
*result = thrust::stable_partition(exec, first, last, pred);
#else
*is_supported = false;
#endif
}


Expand Down Expand Up @@ -359,12 +355,8 @@ template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typen
__global__
void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 stencil_first, Predicate pred, Iterator3 result, Iterator4 is_supported)
{
#if (__CUDA_ARCH__ >= 200)
*is_supported = true;
*result = thrust::stable_partition(exec, first, last, stencil_first, pred);
#else
*is_supported = false;
#endif
}


Expand Down
4 changes: 0 additions & 4 deletions testing/cuda/sort.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,12 +8,8 @@ template<typename ExecutionPolicy, typename Iterator, typename Compare, typename
__global__
void sort_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Compare comp, Iterator2 is_supported)
{
#if (__CUDA_ARCH__ >= 200)
*is_supported = true;
thrust::sort(exec, first, last, comp);
#else
*is_supported = false;
#endif
}


Expand Down
4 changes: 0 additions & 4 deletions testing/cuda/sort_by_key.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,12 +8,8 @@ template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typen
__global__
void sort_by_key_kernel(ExecutionPolicy exec, Iterator1 keys_first, Iterator1 keys_last, Iterator2 values_first, Compare comp, Iterator3 is_supported)
{
#if (__CUDA_ARCH__ >= 200)
*is_supported = true;
thrust::sort_by_key(exec, keys_first, keys_last, values_first, comp);
#else
*is_supported = false;
#endif
}


Expand Down
8 changes: 0 additions & 8 deletions thrust/detail/type_traits.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -568,15 +568,7 @@ template<typename T>

struct largest_available_float
{
#if defined(__CUDA_ARCH__)
# if (__CUDA_ARCH__ < 130)
typedef float type;
# else
typedef double type;
# endif
#else
typedef double type;
#endif
};

// T1 wins if they are both the same size
Expand Down
2 changes: 1 addition & 1 deletion thrust/system/cuda/config.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,7 @@
#include <cub/util_namespace.cuh>

#if defined(__CUDACC__) || defined(_NVHPC_CUDA)
# if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__>= 350 && defined(__CUDACC_RDC__))
# if !defined(__CUDA_ARCH__) || defined(__CUDACC_RDC__)
# define __THRUST_HAS_CUDART__ 1
# define THRUST_RUNTIME_FUNCTION __host__ __device__ __forceinline__
# else
Expand Down