Неверные результаты с CUB ReduceByKey при указании gencode

Question

В одном из моих проектов я вижу некоторые неправильные результаты при использовании CUB's DeviceReduce::ReduceByKey. Тем не менее, использование тех же входов / выходов с thrust::redu_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);

}

Когда я включаю "-gencode arch=compute_35,code=sm_35" (для Kepler GTX Titan), это дает неправильные результаты, но когда я полностью опускаю эти флаги, это работает.

$ 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) и увидеть то же самое поведение.

Правильное ли решение опустить эти флаги компилятора?

попробовал на одной машине с:

• куда 8.0
• Ubuntu 16.04
• gcc 5.4.0
• детеныш 1.6.4
• Kepler GTX Titan (вычислительная мощность 3.5)

и другой с:

• куда 8.0
• Ubuntu 16.04
• gcc 5.4.0
• детеныш 1.6.4
• Pascal GTX 1080 Ti (вычислительная мощность 6.1)

Answer 1

Похоже, вы должны подать отчет об ошибке на странице проблем хранилища CUB.

Изменить: я могу воспроизвести эту проблему:

[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 Mint 18.1
• Ядро Linux: 4.4.0
• GCC: 5.4.0-6ubuntu1 ~ 16.04.4
• детеныш: 1.6.4
• Графический процессор: GTX 650 Ti (Compute Capability 3.0)