使用CUDA

Question

查找未知尺寸的区域的最大值在Array

说我有一个包含所有不同数目[45,21,764,234,7,0,12,55,...]

然后，我有另一个数组B[4000]与表示在阵列A区域的位置值的阵列A[4000]如果它是区域的一部分，则为1，如果不是，则为0。如果1's彼此相邻，则表示它们是同一区域的一部分，如果它们彼此不相邻（1's之间有0），则它们是不同区域的一部分。

ex。 B = [1,1,1,0,1,1,0,0...]就是说我想找到在first three numbers in array A的区域中的最大值，并且在5th and 6th numbers in array A, etc. 所以，我可以产生保持在每个由B表示的区域中的A最大值的数组C[4000]的最大数量，和一个0在不是地区的一部分的区域。

因此，在这种情况下C = [764,764,764,0,7,7,0,0...]

可以有从0 to 2,000 regions在任何地方，和区域的长度可以从2 to 4,000 numbers long范围。我从来不知道有多少区域或区域的大小不同。

我一直在想出一个可以实现这个结果的CUDA内核。它需要尽可能快地完成，因为它实际上将被用于图像，这只是一个简单的例子。我所有的想法，比如使用简化，只在只有一个区域跨越所有4000数字A时才起作用。然而，我不认为我可以在这里使用缩减，因为可能有多个区域由1到3996之间的空格（0's）分隔，阵列中的缩减会导致我分散区域的跟踪。或者，内核有太多的循环，如果在它的语句要快如

int intR = 0; 
while(B[blockIdx.x * blockDim.x + threadIdx.x + intR] > 0){ 
    intMaxR = intMaxR < A[blockIdx.x * blockDim.x + threadIdx.x + intR] ? A[blockIdx.x * blockDim.x + threadIdx.x + intR] : intMaxR; 
    intR++; 
} 

int intL = 0; 
while(B[blockIdx.x * blockDim.x + threadIdx.x - intL] > 0){ 
    intMaxL = intMaxL < A[blockIdx.x * blockDim.x + threadIdx.x - intL] ? A[blockIdx.x * blockDim.x + threadIdx.x + intL] : intMaxL; 
    intL++; 
} 

intMax = intMaxR > intMaxL ? intMaxR : intMaxL; 

for(int i = 0; i < intR; i++){ 
    C[blockIdx.x * blockDim.x + threadIdx.x + i] = intMax; 
} 
for(int i = 0; i < intL; i++){ 
    C[blockIdx.x * blockDim.x + threadIdx.x - i] = intMax; 
} 

显然代码甚至共享内存慢，是不是真正得到了CUDA的并行性质的优势。有没有人有任何想法如何或如果这可以在CUDA中有效地完成？

在此先感谢。

Answer 1

一种可能的方法是使用thrust。

一个可能的序列将是这样的：

1. 使用 thrust::reduce_by_key 以产生用于每个范围的最大值。
2. 使用thrust :: adjacent_difference描绘每个范围的起始点
3. 对步骤2的结果使用包含式扫描来生成聚集索引，即将用于选择缩小值的索引（来自步骤1）将会出现在输出向量的每个位置。
4. 使用thrust::gather_if使用步骤3中生成的聚集索引，选择性地将缩小的值放置到输出向量中的适当位置（B向量中存在1）。

这里的一个充分的工作代码演示此，使用像您的示例A和B载体：

关于示例

#include <iostream> 
#include <thrust/device_vector.h> 
#include <thrust/adjacent_difference.h> 
#include <thrust/reduce.h> 
#include <thrust/copy.h> 
#include <thrust/transform_scan.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/functional.h> 

#define DSIZE 8 

template <typename T> 
struct abs_val : public thrust::unary_function<T, T> 
{ 
    __host__ __device__ 
    T operator()(const T& x) const 
    { 
    if (x<0) return -x; 
    else return x; 
    } 
}; 

template <typename T> 
struct subtr : public thrust::unary_function<T, T> 
{ 
    const T val; 
    subtr(T _val): val(_val) {} 
    __host__ __device__ 
    T operator()(const T& x) const 
    { 
    return x-val; 
    } 
}; 

int main(){ 

    int A[DSIZE] = {45,21,764,234,7,0,12,55}; 
    int B[DSIZE] = {1,1,1,0,1,1,0,0}; 
    thrust::device_vector<int> dA(A, A+DSIZE); 
    thrust::device_vector<int> dB(B, B+DSIZE); 
    thrust::device_vector<int> dRed(DSIZE); 
    thrust::device_vector<int> diffB(DSIZE); 
    thrust::device_vector<int> dRes(DSIZE); 

    thrust::reduce_by_key(dB.begin(), dB.end(), dA.begin(), thrust::make_discard_iterator(), dRed.begin(), thrust::equal_to<int>(), thrust::maximum<int>()); 
    thrust::adjacent_difference(dB.begin(), dB.end(), diffB.begin()); 
    thrust::transform_inclusive_scan(diffB.begin(), diffB.end(), diffB.begin(), abs_val<int>(), thrust::plus<int>()); 
    thrust::gather_if(thrust::make_transform_iterator(diffB.begin(), subtr<int>(B[0])), thrust::make_transform_iterator(diffB.end(), subtr<int>(B[0])), dB.begin(), dRed.begin(), dRes.begin()); 
    thrust::copy(dRes.begin(), dRes.end(), std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 
    return 0; 
} 

注：

1. reduce_by_key正在产生降低的值（最大值），用于每个 连续的0序列或 B中的1个序列。您只需真正需要 这1个序列的最大值。我们将通过gather_if函数丢弃最大值为0的序列 。
2. 我允许该B载体可以与任一 1个序列或0的序列开始，通过使用transform_iterator 处理步骤2的矢量结果，减去从每个B载体的第一 值的可能性收集索引。
3. adjacent_difference操作将产生1或-1到 划定新序列的开始。为了扫描目的（即生成聚集索引），我使用带有abs_val函子的变换包含变换函数variant_inclusive_scan来平等对待它们。

4. 上面的代码应该产生的结果符合期望C输出向量，就像这样：

   $ nvcc -arch=sm_20 -o t53 t53.cu 
   $ ./t53 
   764 764 764 0 7 7 0 0 
   $ 
   

我们可以使用thrust::placeholders进一步简化上面的代码，省去了多余的仿函数定义的需要：

#include <iostream> 
#include <thrust/device_vector.h> 
#include <thrust/adjacent_difference.h> 
#include <thrust/reduce.h> 
#include <thrust/copy.h> 
#include <thrust/transform_scan.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/functional.h> 

#define DSIZE 2000000 
using namespace thrust::placeholders; 

typedef int mytype; 

int main(){ 

    mytype *A = (mytype *)malloc(DSIZE*sizeof(mytype)); 
    int *B = (int *)malloc(DSIZE*sizeof(int)); 
    for (int i = 0; i < DSIZE; i++){ 
    A[i] = (rand()/(float)RAND_MAX)*10.0f; 
    B[i] = rand()%2;} 
    thrust::device_vector<mytype> dA(A, A+DSIZE); 
    thrust::device_vector<int> dB(B, B+DSIZE); 
    thrust::device_vector<mytype> dRed(DSIZE); 
    thrust::device_vector<int> diffB(DSIZE); 
    thrust::device_vector<mytype> dRes(DSIZE); 

    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start); 
    thrust::reduce_by_key(dB.begin(), dB.end(), dA.begin(), thrust::make_discard_iterator(), dRed.begin(), thrust::equal_to<mytype>(), thrust::maximum<mytype>()); 
    thrust::adjacent_difference(dB.begin(), dB.end(), diffB.begin()); 
    thrust::transform_inclusive_scan(diffB.begin(), diffB.end(), diffB.begin(), _1*_1, thrust::plus<int>()); 
    thrust::gather_if(thrust::make_transform_iterator(diffB.begin(), _1 - B[0]), thrust::make_transform_iterator(diffB.end(), _1 - B[0]), dB.begin(), dRed.begin(), dRes.begin()); 
    cudaEventRecord(stop); 
    cudaEventSynchronize(stop); 
    float et; 
    cudaEventElapsedTime(&et, start, stop); 
    std::cout<< "elapsed time: " << et << "ms " << std::endl; 
    thrust::copy(dRes.begin(), dRes.begin()+10, std::ostream_iterator<mytype>(std::cout, " ")); 
    std::cout << std::endl; 
    return 0; 
} 

（I已经修改了上述占位符代码还包括生成一个较大尺寸的数据集的，以及一些基本定时APPA鼠标）