Résultats incorrects avec CUB ReduceByKey lors de la spécification de gencode

Question

Dans l'un de mes projets, des résultats incorrects s'affichent lors de l'utilisation de CUB de DeviceReduce :: ReduceByKey. Cependant, utiliser les mêmes entrées/sorties avec thrust :: reduce_by_key produit les résultats attendus.

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

} 

Quand j'Include "arc -gencode = compute_35, code = sm_35" (pour Kepler GTX Titan), il produit des résultats erronés, mais quand je laisse ces drapeaux entièrement, cela fonctionne. J'utilise une poignée d'autres appels CUB sans problème, juste celui-ci se comporte mal. J'ai également essayé d'exécuter ce code sur un GTX 1080 Ti (avec compute_61, sm_61) et voir le même comportement.

La bonne solution pour omettre ces drapeaux de compilation?

essayé sur une machine avec:

• cuda 8,0
• ubuntu 16.04
• gcc 5.4.0
• Cub 1.6.4
• Kepler GTX Titan (capacité de calcul 3.5)

et un autre avec:

• cuda 8,0
• ubuntu 16.04
• gcc 5.4.0
• Cub 1.6.4
• Pascal GTX 1080 Ti (Compute capacité 6,1)

Answer 1

On dirait que vous devez déposer une rapport de bug au CUB repository issues page.

Edit: Je peux reproduire ce problème:

[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 

Infodesks:

• CUDA: 8.0.61
• pilote NVidia: 375,39
• Distribution: GNU/Linux Menthe 18.1
• Noyau Linux: 4.4.0
• GCC: 5.4.0-6ubuntu1 ~ 16.04.4
• Cub: 1.6.4
• GPU: GTX 650 Ti (capacité Compute 3,0)