多k-selection的cuda和cub实现答案

【问题标题】：cuda and cub implementation of multiple k-selection多k-selection的cuda和cub实现
【发布时间】：2015-08-12 18:47:40
【问题描述】：

我正在尝试并行实现多个 top-k 选择，其中每个选择从 n 个元素的列表中选择 k 个元素，并且有 m 个这样的任务要并行执行。我用 cub 来做到这一点。我遇到了一个奇怪的错误，我不知道我在哪里做错了。我觉得我的理解可能有一些明显的错误，有人可以帮我检查一下吗？

编辑：

我通过在包含free() 的两个代码部分中的每一个之前添加两个cudaDeviceSynchronize() 调用来使其工作。所以现在我的问题是，free 的行为是否与cudaFree 不同，因为不允许立即进行异步调用，而不是我在这里提出的另一个问题：Does cudaFree after asynchronous call work?

// 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);
}

【问题讨论】：

标签： cuda cub

【解决方案1】：

虽然我能够让它工作，但这个实现比单线程 CPU 代码执行得慢。我最终用堆排序实现了这个并将堆放在共享内存中。性能不错。

【讨论】：