0

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

编辑：

我通过cudaDeviceSynchronize()在包含free(). 所以现在我的问题是，does 的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);
}
4

1 回答 1

1

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

于 2014-01-19T04:43:26.817 回答