GUBodeを指定すると、CUB ReduceByKeyの結果が正しくない

Question

私のプロジェクトの1つでは、CUBの DeviceReduce :: ReduceByKeyを使用すると誤った結果が表示されます。しかし、thrust :: reduce_by_keyと同じ入出力を使用すると、期待される結果が得られます。

#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); 

} 

私は（ケプラーGTXタイタン用）「-gencodeアーチ= compute_35、コード= sm_35」が含まれる場合は、それは間違った結果を生成しますが、私は完全にこれらのフラグを去るとき、それは動作します。私は他のCUBの一握りは、この1つだけが不正な動作され、問題なく呼び出しを使用

$ 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 

。また、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
• パスカルGTX 1080あなたが提出すべきのようなチタン（計算機能6.1）

Answer 1

が鳴りますバグ報告はCUB repository issues pageです。

編集：

[joeuser@myhost:/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). 
[joeuser@myhost:/tmp]$ ./a 
0, 4000 
1, 1000 
2, 2680 
[joeuser@myhost:/tmp]$ nvcc -I/opt/cub -o a a.cu -gencode arch=compute_30,code=sm_30 
[joeuser@myhost:/tmp]$ ./a 
0, 4000 
1, 1000 
2, 512 

関連情報：

• CUDA：8.0.61
• のnVIDIAドライバ：375.39
• 配布：GNU/Linuxの私は、この問題を再現することができますミント18.1
• Linuxカーネル：4.4.0
• GCC：5.4.0-6ubuntu1〜16.04.4
• カブ：1.6.4
• GPU：GTX 650のTi（計算能力3.0）