cuda и cub реализация множественного k-выделения

Question

Я пытаюсь реализовать множественный выбор top-k параллельно, где каждый выбор выбирает k элементов из списка из n элементов, и m таких задач должны выполняться параллельно. Я использую детеныша, чтобы сделать это. Я получил странную ошибку, и я не знаю, где я сделал не так. Я чувствую, что, возможно, допустил некоторую очевидную ошибку в моем понимании, может ли кто-нибудь помочь мне проверить?

РЕДАКТИРОВАТЬ:

Я сделал это, добавив два cudaDeviceSynchronize() звонки, перед каждым из двух разделов кода, который содержит free(), Итак, теперь мой вопрос: free ведет себя иначе, чем cudaFreeв том, что асинхронные вызовы, выполняющиеся немедленно, запрещены, в отличие от другого вопроса, который я задал здесь: работает ли cudaFree после асинхронного вызова?

// Assume dtop has size k x m and dmat has size n x m, where k < n
// Each column of dtop is supposed to obtain the top-k indices of 
// elements from the corresponding column in dmat.
template<typename ValueType, typename IndexType>
void TopKPerColumn_cub_test(DenseMatrix<IndexType, MemDev> dtop,
    DenseMatrix<ValueType, MemDev, Const> dmat);

template<typename T>
struct SelectLE {
  T x_;
  __device__ SelectLE(const T& x):x_(x){}
  __device__ bool operator() (const T& a) {
    return a > x_;
  }
};

template<typename ValueType, typename IndexType>
__global__ void k_TopKPerColumn_cub_test(DenseMatrix<IndexType, MemDev> dtop,
    DenseMatrix<ValueType, MemDev, Const> dmat) {
  int n = dmat.num_rows();
  int k = dtop.num_rows();

  cub::DoubleBuffer<ValueType> keys;
  keys.d_buffers[0] = reinterpret_cast<ValueType*>(
      malloc(sizeof(ValueType) * n));
  keys.d_buffers[1] = reinterpret_cast<ValueType*>(
      malloc(sizeof(ValueType) * n));
  memcpy(keys.d_buffers[keys.selector], dmat.get_col(blockIdx.x).data(),
      sizeof(ValueType) * n);

  void* temp_storage = 0;
  size_t temp_storage_size = 0;
  cub::DeviceRadixSort::SortKeysDescending(
      temp_storage, temp_storage_size, keys, n);
  temp_storage = malloc(temp_storage_size);
  cub::DeviceRadixSort::SortKeysDescending(
      temp_storage, temp_storage_size, keys, n);
  ValueType kth = keys.Current()[k-1];

  free(temp_storage);
  free(keys.d_buffers[0]);
  free(keys.d_buffers[1]);

  temp_storage = 0;
  temp_storage_size = 0;
  int* nb_selected = reinterpret_cast<int*>(malloc(sizeof(int)));
  SelectLE<ValueType> selector(kth);

  cub::DeviceSelect::If(temp_storage, temp_storage_size,
      const_cast<ValueType*>(dmat.get_col(blockIdx.x).data()),
      dtop.get_col(blockIdx.x).data(),
      nb_selected, n, selector);
  temp_storage = malloc(temp_storage_size);
  cub::DeviceSelect::If(temp_storage, temp_storage_size,
      const_cast<ValueType*>(dmat.get_col(blockIdx.x).data()),
      dtop.get_col(blockIdx.x).data(),
      nb_selected, n, selector);

  free(nb_selected);
  free(temp_storage);
}

template<typename ValueType, typename IndexType>
void TopKPerColumn_cub_test(DenseMatrix<IndexType, MemDev> dtop,
    DenseMatrix<ValueType, MemDev, Const> dmat) {
  k_TopKPerColumn_cub_test<<<dtop.num_cols(), 1>>>(dtop, dmat);
}

Answer 1

Хотя я могу заставить его работать, эта реализация работает медленнее, чем однопоточный код процессора. В конце концов я реализовал это с помощью сортировки кучи и поместил кучу в общую память. Производительность хорошая.