如何使用 CUDA C 快速压缩稀疏数组?

Posted

技术标签:

【中文标题】如何使用 CUDA C 快速压缩稀疏数组?【英文标题】:How to quickly compact a sparse array with CUDA C? 【发布时间】:2013-01-10 12:41:41 【问题描述】:

总结

设备内存中的数组 [A - B - - - C] 但想要 [A B C] - 使用 CUDA C 最快的方法是什么?

上下文

我在设备 (GPU) 内存上有一个整数数组 A。在每次迭代中,我随机选择几个大于 0 的元素并从中减去 1。我维护了一个排序的查找数组L,其中包含等于 0 的元素:

Array A:
       @ iteration i: [0 1 0 3 3 2 0 1 2 3]
   @ iteration i + 1: [0 0 0 3 2 2 0 1 2 3]

Lookup for 0-elements L:
       @ iteration i: [0 - 2 - - - 6 - - -]  ->  want compacted form: [0 2 6]
   @ iteration i + 1: [0 1 2 - - - 6 - - -]  ->  want compacted form: [0 1 2 6]

(这里,我随机选择元素14 减去1。在我在CUDA C 中的实现中,每个线程映射到A 中的一个元素,因此查找数组是稀疏的防止数据竞争并保持有序的顺序(例如[0 1 2 6] 而不是[0 2 6 1])。

稍后,我将只对那些等于 0 的元素进行一些操作。因此我需要压缩我的稀疏查找数组L,以便我可以将线程映射到 0 元素。

因此,使用 CUDA C 在设备内存上压缩稀疏数组的最有效方法是什么?

非常感谢。

【问题讨论】:

您可以考虑使用thrust stream compaction。 谢谢 - 标准 CUDA 安装附带推力吗?由于我不是系统管理员,如果库可用,我如何检查 Unix 机器?谢谢。 是的,假设是最新版本的 CUDA。如果你有一个像/usr/local/cuda/include/thrust 这样的目录,那么你就有了推力。 Thrust 完全是模板化/包含的代码,因此无需担心普通的库。您可能对quick start guide 感兴趣。 感谢@RobertCrovella,但我看不到 C 用户的任何示例用法 - 只有我不熟悉的 C++。例如,您如何在 CUDA C 设备内存中的数组上调用 thrust::copy_if() cuSPARSE 库提供cusparseSdense2csr() 将矩阵从密集格式转换为稀疏格式。它应该非常高效,但可能效率低于thrust::copy_if 【参考方案1】:

假设我有:

int V[] = 1, 2, 0, 0, 5;

而我想要的结果是:

int R[] = 1, 2, 5

实际上,我们正在删除零元素,或者仅复制非零元素。

#include <thrust/device_ptr.h>
#include <thrust/copy.h>
#include <stdio.h>
#define SIZE 5

#define cudaCheckErrors(msg) \
    do  \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess)  \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
         \
     while (0)

  struct is_not_zero
  
    __host__ __device__
    bool operator()(const int x)
    
      return (x != 0);
    
  ;



int main()

  int V[] = 1, 2, 0, 0, 5;
  int R[] = 0, 0, 0, 0, 0;
  int *d_V, *d_R;

  cudaMalloc((void **)&d_V, SIZE*sizeof(int));
  cudaCheckErrors("cudaMalloc1 fail");
  cudaMalloc((void **)&d_R, SIZE*sizeof(int));
  cudaCheckErrors("cudaMalloc2 fail");

  cudaMemcpy(d_V, V, SIZE*sizeof(int), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy1 fail");

  thrust::device_ptr<int> dp_V(d_V);
  thrust::device_ptr<int> dp_R(d_R);
  thrust::copy_if(dp_V, dp_V + SIZE, dp_R, is_not_zero());

  cudaMemcpy(R, d_R, SIZE*sizeof(int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy2 fail");

  for (int i = 0; i<3; i++)
    printf("R[%d]: %d\n", i, R[i]);

  return 0;



结构定义为我们提供了一个测试零元素的函子。请注意,在推力中,没有内核,我们也没有直接编写设备代码。这一切都发生在幕后。我绝对建议您熟悉quick start guide,以免把这个问题变成推力教程。

在查看了 cmets 之后,我认为这个修改后的代码版本可以解决 cuda 4.0 的问题:

#include <thrust/device_ptr.h>
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <stdio.h>
#define SIZE 5

  struct is_not_zero
  
    __host__ __device__
    bool operator()(const int x)
    
      return (x != 0);
    
  ;



int main()

  int V[] = 1, 2, 0, 0, 5;
  int R[] = 0, 0, 0, 0, 0;

  thrust::host_vector<int> h_V(V, V+SIZE);
  thrust::device_vector<int> d_V = h_V;
  thrust::device_vector<int> d_R(SIZE, 0);

  thrust::copy_if(d_V.begin(), d_V.end(), d_R.begin(), is_not_zero());
  thrust::host_vector<int> h_R = d_R;

  thrust::copy(h_R.begin(), h_R.end(), R);

  for (int i = 0; i<3; i++)
    printf("R[%d]: %d\n", i, R[i]);

  return 0;



【讨论】:

谢谢。我尝试了您的解决方案,但出现编译错误:[...]/cuda/4.0.17/cuda/bin/../include/thrust/detail/device/cuda/copy_i‌​f.inl(71): error: more than one instance of overloaded function "min" matches the argument list: function "min(int, int)" function "min(unsigned int, unsigned int)" [...] argument types are: (long, const long) detected during: instantiation of "void thrust::detail::device::cuda::reduce_intervals&lt;CTA_SIZE,InputIterator,IndexType,‌​OutputIterator,BinaryFunction&gt;(InputIterator, IndexType, IndexType, OutputIterator, BinaryFunction) [...] 其中,[...] 是截断。 你拿了我发布的代码,并试图编译它?或者您是否进行了任何更改或添加?看起来您正在使用 CUDA 4.0。我在 cuda 4.2 和 cuda 5.0 上测试过,但不是 4.0 谢谢罗伯特,我完全按照您发布的代码编译了代码。知道为什么 CUDA 4.0 会抱怨吗? cuda 4 已经很老了。现在2岁多了。尝试将 -m32 添加到您的 nvcc 编译命令行。 干杯。现在我得到错误:In file included from /usr/include/features.h:371, from [...]/cuda/4.0.17/cuda/bin/../include/host_config.h:114, from [...]/cuda/4.0.17/cuda/bin/../include/cuda_runtime.h:59, from &lt;command-line&gt;:0: /usr/include/gnu/stubs.h:7:27: error: gnu/stubs-32.h: No such file or directory。感谢您的耐心等待。

以上是关于如何使用 CUDA C 快速压缩稀疏数组?的主要内容,如果未能解决你的问题,请参考以下文章

热榜!!!数据结构与算法:C语言版---数组与稀疏矩阵---强势来袭!

热榜!!!数据结构与算法:C语言版---数组与稀疏矩阵---强势来袭!

如何对存储为“压缩稀疏行”的矩阵进行稀疏矩阵索引?

用C试一下稀疏矩阵的快速转置

一文带你读懂非结构化稀疏模型压缩和推理优化技术

一文带你读懂非结构化稀疏模型压缩和推理优化技术