c++ - 定义自定义扫描运算符

我正在尝试在cub中创建自己的扫描运算符。它现在可以工作，但仅适用于小于1024的数组，这让我认为它仅适用于块。这是我的代码:

#include "cub/cub.cuh"
using namespace cub;


typedef int mytype;

struct CustomMin
{
    template <typename T>
    __host__ __device__
    CUB_RUNTIME_FUNCTION __forceinline__
    mytype operator()(const T &a, const T &b) const {
        return (b < a) ? b : a;
    }
};


int main(int argc, char *argv[])
{

    int num_items = 512;
    mytype *h_in;
    mytype *h_out;
    CustomMin    min_op;
    const size_t size = num_items * sizeof(mytype);
    h_in = (mytype*)malloc(size);
    h_out = (mytype*)malloc(size);
    mytype *d_in = NULL;
    cudaMalloc(&d_in, size);
    mytype *d_out = NULL;
    cudaMalloc(&d_out, size);
    for (int i = 0; i < num_items; i++) {
        h_in[i] = i;
    }
    cudaMemcpy(d_in, h_in, size, cudaMemcpyHostToDevice);
    void *d_temp_storage = NULL;
    size_t temp_storage_bytes = 0;
    DeviceScan::InclusiveScan(d_temp_storage, temp_storage_bytes, d_in, d_out, min_op, num_items);
    cudaMalloc(&d_temp_storage, temp_storage_bytes);
    DeviceScan::InclusiveScan(d_temp_storage, temp_storage_bytes, d_in, d_out, min_op, num_items);
    cudaMemcpy(h_out, d_out, size, cudaMemcpyDeviceToHost);
    printf("done!\n");
    return 0;
}

对于较大的输入大小，它总是挂起。

最佳答案

使用CUB 1.4.1，我可以像这样编译时重现挂起:

nvcc -arch=sm_35 -o t25 t25.cu

将发布的代码中的num_items更改为2048之后。

根据我的测试，该问题似乎已在cub 1.5.1中修复。请更新到最新的CUB版本。