Cuda 向量类型的推力支持
Thrust Support of Cuda Vector Types
我目前正在尝试使用 thrust::upper_bound 功能。我 运行 对我提供给函数的参数有疑问。我想使用 CUDA 矢量类型，特别是 double3，但是当我使用这种类型时，我遇到了几个推力库错误。
我是运行的代码块如下：
/********************************************************************************
eos_search_gpu()
purpose --- kernel to find the upper bound index for the
interpolation values
arguments --
y --- input double3 values for which we are searching
my --- input int number of values for which we are searching
x --- input double3 array of structs containin the data table
values for x, y, and f corresponding to structs
".x", ".y", and ".z"
n --- input int number of data values in the table
dim_x --- input int number of data values in the x-direcion of table
j[] --- input/output int[] array of int'sthat contains
the index of the (x,y,f) position of the upper bound
library calls --
__host__ __device__ ForwardIterator thrust::upper_bound(
const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
ForwardIterator first,
ForwardIterator last,
const LessThanComparable & value
)
exec --- the execution policy to use for parallelization
first --- the beginning of the ordered sequence
last --- the end of the ordered sequence
value --- the value to be searched.
Returns: the furthermost iterator i, such that value < *i is false
const detail::seq_t thrust::seq
an execution policy which requires analgorithm invocation to execute
sequentially in the current thread.
********************************************************************************/
__global__ void eos_search_gpu(const double3* y, const int my,
const double3* x, const int n,
const int dim_x, int * j){
int i = threadIdx.x + blockDim.x * blockIdx.x;
if ( i < my) {
const double ptr = thrust::upper_bound(thrust::seq, x[0].y , x[n-1].y, y[i].y);
j[i] = (ptr - x[i].y - 1);
}
}
显示的错误信息如下：
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: a class or namespace qualified name is required
detected during:
instantiation of class "thrust::iterator_traits<T> [with T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]"
Interpolation_cuda.cu(254): here
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: global-scope qualifier (leading "::") is not allowed
detected during:
instantiation of class "thrust::iterator_traits<T> [with T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]"
Interpolation_cuda.cu(254): here
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: expected a ";"
detected during:
instantiation of class "thrust::iterator_traits<T> [with T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]"
Interpolation_cuda.cu(254): here
我想知道 thrust 是否支持使用 CUDA 矢量类型，或者我只是做错了什么。
您需要满足推力算法的所有预期输入类型。你没有这样做，因为你定义的几乎每个数量都不符合预期的推力。
首先，我们需要实际的迭代器。在设备代码中，这意味着指针。 Thrust 需要能够取消引用 iterator/pointer，然后您必须指示 thrust 如何处理该数量。为此，我们需要一个适当定义的函子。您可能希望阅读 upper_bound 的 thrust quick start guide to understand functor definition and usage. Finally, the sensible pointer/iterator here refers to a double3 type, so we'll need to craft nearly everything to work with double3. Note that we need to choose the version，它允许定义我们自己的自定义仿函数，因此我们可以操纵 double3 数量（当我们取消引用 iterators/pointers) 正确。
这可能有帮助：
#include <thrust/binary_search.h>
#include <thrust/execution_policy.h>
struct my_comp_functor{
template <typename T>
__host__ __device__
bool operator()(T &t1, T &t2) {
return (t1.y < t2.y);}
};
__global__ void eos_search_gpu(const double3* y, const int my,
const double3* x, const int n,
const int dim_x, int * j, my_comp_functor my_comp){
int i = threadIdx.x + blockDim.x * blockIdx.x;
if ( i < my) {
const double3 *ptr = thrust::upper_bound(thrust::seq, x, x+n, y[i], my_comp);
j[i] = (ptr[0].y - x[i].y - 1);
}
}
int main(){
double3 *d_y, *d_x;
int *d_j;
cudaMalloc(&d_y, 1024);
cudaMalloc(&d_x, 1024);
cudaMalloc(&d_j, 1024);
struct my_comp_functor my_obj;
eos_search_gpu<<<1,1>>>(d_y, 0, d_x, 0, 0, d_j, my_obj);
cudaDeviceSynchronize();
}
（上面的代码在CUDA 9.2上对我来说没有编译错误，但显然不是设计成functional/useful）
最后，我觉得你把一个 double 数量塞进 j[i]（一个整数）但它是你的代码。
此外，我可能在那个仿函数中弄错了顺序，所以您可能需要将 < 更改为 >。
调用这个内核的时候，注意我加了一个参数；您需要在主机代码中实例化一个 my_comp_functor 对象，然后将其传递给适当位置的内核。
最后，您似乎在进行矢量化搜索，请注意推力 vectorized searches available 可能不需要此内核。
我目前正在尝试使用 thrust::upper_bound 功能。我 运行 对我提供给函数的参数有疑问。我想使用 CUDA 矢量类型，特别是 double3，但是当我使用这种类型时，我遇到了几个推力库错误。
我是运行的代码块如下：
/********************************************************************************
eos_search_gpu()
purpose --- kernel to find the upper bound index for the
interpolation values
arguments --
y --- input double3 values for which we are searching
my --- input int number of values for which we are searching
x --- input double3 array of structs containin the data table
values for x, y, and f corresponding to structs
".x", ".y", and ".z"
n --- input int number of data values in the table
dim_x --- input int number of data values in the x-direcion of table
j[] --- input/output int[] array of int'sthat contains
the index of the (x,y,f) position of the upper bound
library calls --
__host__ __device__ ForwardIterator thrust::upper_bound(
const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
ForwardIterator first,
ForwardIterator last,
const LessThanComparable & value
)
exec --- the execution policy to use for parallelization
first --- the beginning of the ordered sequence
last --- the end of the ordered sequence
value --- the value to be searched.
Returns: the furthermost iterator i, such that value < *i is false
const detail::seq_t thrust::seq
an execution policy which requires analgorithm invocation to execute
sequentially in the current thread.
********************************************************************************/
__global__ void eos_search_gpu(const double3* y, const int my,
const double3* x, const int n,
const int dim_x, int * j){
int i = threadIdx.x + blockDim.x * blockIdx.x;
if ( i < my) {
const double ptr = thrust::upper_bound(thrust::seq, x[0].y , x[n-1].y, y[i].y);
j[i] = (ptr - x[i].y - 1);
}
}
显示的错误信息如下：
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: a class or namespace qualified name is required
detected during:
instantiation of class "thrust::iterator_traits<T> [with T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]"
Interpolation_cuda.cu(254): here
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: global-scope qualifier (leading "::") is not allowed
detected during:
instantiation of class "thrust::iterator_traits<T> [with T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]"
Interpolation_cuda.cu(254): here
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_traits.h(45): error: expected a ";"
detected during:
instantiation of class "thrust::iterator_traits<T> [with T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/iterator/detail/iterator_traits.inl(53): here
instantiation of class "thrust::iterator_difference<Iterator> [with Iterator=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/sequential/binary_search.h(102): here
instantiation of "ForwardIterator thrust::system::detail::sequential::upper_bound(thrust::system::detail::sequential::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(83): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &, StrictWeakOrdering) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double, StrictWeakOrdering=thrust::system::detail::generic::detail::binary_search_less]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/binary_search.inl(225): here
instantiation of "ForwardIterator thrust::system::detail::generic::upper_bound(thrust::execution_policy<DerivedPolicy> &, ForwardIterator, ForwardIterator, const T &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, T=double]"
/opt/cudatoolkit/9.1/bin/../targets/x86_64-linux/include/thrust/detail/binary_search.inl(69): here
instantiation of "ForwardIterator thrust::upper_bound(const thrust::detail::execution_policy_base<DerivedPolicy> &, ForwardIterator, ForwardIterator, const LessThanComparable &) [with DerivedPolicy=thrust::detail::seq_t, ForwardIterator=double, LessThanComparable=double]"
Interpolation_cuda.cu(254): here
我想知道 thrust 是否支持使用 CUDA 矢量类型，或者我只是做错了什么。
您需要满足推力算法的所有预期输入类型。你没有这样做，因为你定义的几乎每个数量都不符合预期的推力。
首先，我们需要实际的迭代器。在设备代码中，这意味着指针。 Thrust 需要能够取消引用 iterator/pointer，然后您必须指示 thrust 如何处理该数量。为此，我们需要一个适当定义的函子。您可能希望阅读 upper_bound 的 thrust quick start guide to understand functor definition and usage. Finally, the sensible pointer/iterator here refers to a double3 type, so we'll need to craft nearly everything to work with double3. Note that we need to choose the version，它允许定义我们自己的自定义仿函数，因此我们可以操纵 double3 数量（当我们取消引用 iterators/pointers) 正确。
这可能有帮助：
#include <thrust/binary_search.h>
#include <thrust/execution_policy.h>
struct my_comp_functor{
template <typename T>
__host__ __device__
bool operator()(T &t1, T &t2) {
return (t1.y < t2.y);}
};
__global__ void eos_search_gpu(const double3* y, const int my,
const double3* x, const int n,
const int dim_x, int * j, my_comp_functor my_comp){
int i = threadIdx.x + blockDim.x * blockIdx.x;
if ( i < my) {
const double3 *ptr = thrust::upper_bound(thrust::seq, x, x+n, y[i], my_comp);
j[i] = (ptr[0].y - x[i].y - 1);
}
}
int main(){
double3 *d_y, *d_x;
int *d_j;
cudaMalloc(&d_y, 1024);
cudaMalloc(&d_x, 1024);
cudaMalloc(&d_j, 1024);
struct my_comp_functor my_obj;
eos_search_gpu<<<1,1>>>(d_y, 0, d_x, 0, 0, d_j, my_obj);
cudaDeviceSynchronize();
}
（上面的代码在CUDA 9.2上对我来说没有编译错误，但显然不是设计成functional/useful）
最后，我觉得你把一个 double 数量塞进 j[i]（一个整数）但它是你的代码。
此外，我可能在那个仿函数中弄错了顺序，所以您可能需要将 < 更改为 >。
调用这个内核的时候，注意我加了一个参数；您需要在主机代码中实例化一个 my_comp_functor 对象，然后将其传递给适当位置的内核。
最后，您似乎在进行矢量化搜索，请注意推力 vectorized searches available 可能不需要此内核。