使用CUB ReduceByKey指定gencode时错误的结果

在我的其中一个项目中，使用CUB的 DeviceReduce :: ReduceByKey时看到一些不正确的结果。但是，在thrust :: reduce_by_key中使用相同的输入/输出会产生预期结果。使用CUB ReduceByKey指定gencode时错误的结果

#include "cub/cub.cuh" 

#include <vector> 
#include <iostream> 

#include <cuda.h> 

struct AddFunctor { 
    __host__ __device__ __forceinline__ 
    float operator()(const float & a, const float & b) const { 
    return a + b; 
    } 
} reduction_op; 

int main() { 

    int n = 7680; 

    std::vector <uint64_t> keys_h(n); 
    for (int i = 0; i < 4000; i++) keys_h[i] = 1; 
    for (int i = 4000; i < 5000; i++) keys_h[i] = 2; 
    for (int i = 5000; i < 7680; i++) keys_h[i] = 3; 

    uint64_t * keys; 
    cudaMalloc(&keys, sizeof(uint64_t) * n); 
    cudaMemcpy(keys, &keys_h[0], sizeof(uint64_t) * n, cudaMemcpyDefault); 

    uint64_t * unique_keys; 
    cudaMalloc(&unique_keys, sizeof(uint64_t) * n); 

    std::vector <float> values_h(n); 
    for (int i = 0; i < n; i++) values_h[i] = 1.0; 

    float * values; 
    cudaMalloc(&values, sizeof(float) * n); 
    cudaMemcpy(values, &values_h[0], sizeof(float) * n, cudaMemcpyDefault); 

    float * aggregates; 
    cudaMalloc(&aggregates, sizeof(float) * n); 

    int * remaining; 
    cudaMalloc(&remaining, sizeof(int)); 

    size_t size = 0; 
    void * buffer = NULL; 

    cub::DeviceReduce::ReduceByKey(
    buffer, 
    size, 
    keys, 
    unique_keys, 
    values, 
    aggregates, 
    remaining, 
    reduction_op, 
    n); 

    cudaMalloc(&buffer, sizeof(char) * size); 

    cub::DeviceReduce::ReduceByKey(
    buffer, 
    size, 
    keys, 
    unique_keys, 
    values, 
    aggregates, 
    remaining, 
    reduction_op, 
    n); 

    int remaining_h; 
    cudaMemcpy(&remaining_h, remaining, sizeof(int), cudaMemcpyDefault); 

    std::vector <float> aggregates_h(remaining_h); 
    cudaMemcpy(&aggregates_h[0], aggregates, sizeof(float) * remaining_h, cudaMemcpyDefault); 

    for (int i = 0; i < remaining_h; i++) { 
    std::cout << i << ", " << aggregates_h[i] << std::endl; 
    } 

    cudaFree(buffer); 
    cudaFree(keys); 
    cudaFree(unique_keys); 
    cudaFree(values); 
    cudaFree(aggregates); 
    cudaFree(remaining); 

}

当我有“-gencode ARCH = compute_35，代码= sm_35”（对于开普勒GTX泰坦），它产生错误的结果，但是当我离开这些标志完全出来，它的工作原理。

$ nvcc cub_test.cu 
$ ./a.out 
0, 4000 
1, 1000 
2, 2680 
$ nvcc cub_test.cu -gencode arch=compute_35,code=sm_35 
$ ./a.out 
0, 4000 
1, 1000 
2, 768

我使用了一些其他CUB调用没有问题，只是这一个是行为不端。我也试着在GTX 1080 Ti上运行这个代码（与 compute_61，sm_61）并且看到相同的行为。

是否忽略这些编译器标志是正确的解决方案？

尝试一个机器上：

CUDA 8.0
的ubuntu 16.04
GCC 5.4.0
幼兽1.6.4
开普勒GTX泰坦（计算能力3.5）

和另一个：

CUDA 8.0
的Ubuntu 16.04
GCC 5.4.0
崽1.6.4
帕斯卡尔1080 GTX钛（计算能力6.1）

来源

2017-04-05 smish

我会尝试重现的开普勒这个明天泰坦我们在工作。 – einpoklum

听起来像是你应该提交错误报告在CUB repository issues page。

编辑：我可以重现这个问题：

[[email protected]:/tmp]$ nvcc -I/opt/cub -o a a.cu 
nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning). 
[[email protected]:/tmp]$ ./a 
0, 4000 
1, 1000 
2, 2680 
[[email protected]:/tmp]$ nvcc -I/opt/cub -o a a.cu -gencode arch=compute_30,code=sm_30 
[[email protected]:/tmp]$ ./a 
0, 4000 
1, 1000 
2, 512

使用CUB ReduceByKey指定gencode时错误的结果

回答

相关问题