【问题标题】:Compress "sparse data" with CUDA (CCL: connected component labeling reduction)使用 CUDA 压缩“稀疏数据”(CCL:连接组件标签缩减)
【发布时间】:2019-05-11 02:53:17
【问题描述】:

我有一个 500 万个 32 位整数列表(实际上是 2048 x 2560 图像),其中 90% 是零。非零单元是完全不连续或以任何方式连续的标签(例如 2049、8195、1334300、34320923、4320932)(它是我们自定义连接组件标签 CCL 算法的输出)。我正在使用 NVIDA Tesla K40,所以如果它需要任何前缀扫描工作,我会喜欢它,它使用 SHUFFLE、BALLOT 或任何更高的 CC 功能。

我不需要一个完整的例子,只是一些建议。

为了说明,这里有一篇使用我们的 CCL 算法标记的博客。

其他 blob 将具有不同的唯一标签(例如 13282)。但是所有都将被零包围,并且是椭圆形的。 (我们针对椭圆体优化了 CCL,这就是我们不使用这些库的原因)。但一个副作用是 blob 标签不是连续的数字。我们不关心它们的编号顺序,但我们想要一个标记为 #1 的 blob,另一个标记为 #2,最后一个标记为 #n,其中 n 是图像中 blob 的数量。

标签#1 是什么意思?我的意思是所有 2242 个单元格都应该用 1 替换。所有 13282 个单元格都应该是 #2,等等。

我们 CCL 的最大 blob 数等于 2048x2560。所以我们知道数组的大小。

实际上,罗伯特·克罗维拉(Robert Crovella)在一天前就已经给出了很好的答案。这并不准确,但我现在看到了如何应用答案。所以我不需要更多的帮助。但是他非常慷慨地付出了时间和精力,并要求我用示例重新编写问题,所以我这样做了。

【问题讨论】:

    标签: cuda gpu cudafy.net


    【解决方案1】:

    一种可能的方法是使用以下序列:

    1. thrust::transform - 将输入数据全部转换为1或0:

      0 27 42  0 18 99 94 91  0  -- input data
      0  1  1  0  1  1  1  1  0  -- this will be our "mask vector"
      
    2. thrust::inclusive_scan - 将掩码向量转换为渐进序列:

      0  1  1  0  1  1  1  1  0  -- "mask" vector
      0  1  2  2  3  4  5  6  6  -- "sequence" vector
      
    3. 另一个thrust::transform 用于屏蔽非增加值:

      0  1  1  0  1  1  1  1  0  -- "mask" vector
      0  1  2  2  3  4  5  6  6  -- "sequence" vector
      -------------------------
      0  1  2  0  3  4  5  6  0  -- result of "AND" operation
      

    请注意,我们可以将前两步与thrust::transform_inclusive_scan 结合起来,然后将第三步作为thrust::transform 执行,并使用稍微不同的变换函子。这种修改允许我们省去临时“掩码”向量的创建。

    这是一个完整的示例,展示了使用 thrust::transform_inclusive_scan 的“修改”方法:

    $ cat t635.cu
    #include <iostream>
    #include <stdlib.h>
    
    #include <thrust/device_vector.h>
    #include <thrust/host_vector.h>
    #include <thrust/transform.h>
    #include <thrust/transform_scan.h>
    #include <thrust/generate.h>
    #include <thrust/copy.h>
    
    
    #define DSIZE 20
    #define PCT_ZERO 40
    
    struct my_unary_op
    {
      __host__ __device__
      int operator()(const int data) const
      {
        return (!data) ?  0:1;}
    };
    
    struct my_binary_op
    {
      __host__ __device__
      int operator()(const int d1, const int d2) const
      {
        return (!d1) ? 0:d2;}
    };
    
    int main(){
    
    // generate DSIZE random 32-bit integers, PCT_ZERO% are zero
      thrust::host_vector<int> h_data(DSIZE);
      thrust::generate(h_data.begin(), h_data.end(), rand);
      for (int i = 0; i < DSIZE; i++)
        if ((rand()%100)< PCT_ZERO) h_data[i] = 0;
        else h_data[i] %= 1000;
      thrust::device_vector<int> d_data = h_data;
      thrust::device_vector<int> d_result(DSIZE);
      thrust::transform_inclusive_scan(d_data.begin(), d_data.end(), d_result.begin(), my_unary_op(), thrust::plus<int>());
      thrust::transform(d_data.begin(), d_data.end(), d_result.begin(), d_result.begin(), my_binary_op());
      thrust::copy(d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
      thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
      return 0;
    }
    
    $ nvcc -o t635 t635.cu
    $ ./t635
    0,886,777,0,793,0,386,0,649,0,0,0,0,59,763,926,540,426,0,736,
    0,1,2,0,3,0,4,0,5,0,0,0,0,6,7,8,9,10,0,11,
    $
    

    在我看来,响应更新,这些新信息使问题更难以解决。直方图技术浮现在脑海中,但对 32 位整数(标签)的占用范围没有任何限制,或者对特定标签在数据集中可能重复的次数没有任何限制,直方图技术似乎不切实际。这让我考虑对数据进行排序。

    这样的方法应该可行:

    1. 使用thrust::sort对数据进行排序。
    2. 使用thrust::unique 删除重复项。
    3. 删除重复的排序数据现在为我们提供了输出集 [0,1,2, ...] 的排序。让我们称之为我们的“地图”。我们可以使用parallel binary-search technique 将原始数据集中的每个标签转换为其映射的输出值。

    这个过程对我来说似乎相当“昂贵”。我建议重新考虑上游标记操作,看看是否可以重新设计它以生成更适合高效下游处理的数据集。

    无论如何,这是一个完整的例子:

    $ cat t635.cu
    #include <iostream>
    #include <stdlib.h>
    
    #include <thrust/device_vector.h>
    #include <thrust/host_vector.h>
    #include <thrust/transform.h>
    #include <thrust/generate.h>
    #include <thrust/sort.h>
    #include <thrust/unique.h>
    #include <thrust/copy.h>
    
    
    #define DSIZE 20
    #define PCT_ZERO 40
    #define RNG 10
    
    #define nTPB 256
    
    // sets idx to the index of the first element in a that is
    // equal to or larger than key
    
    __device__ void bsearch_range(const int *a, const int key, const unsigned len_a, unsigned *idx){
      unsigned lower = 0;
      unsigned upper = len_a;
      unsigned midpt;
      while (lower < upper){
        midpt = (lower + upper)>>1;
        if (a[midpt] < key) lower = midpt +1;
        else upper = midpt;
        }
      *idx = lower;
      return;
      }
    
    __global__ void find_my_idx(const int *a, const unsigned len_a,  int *my_data, int *my_idx, const unsigned len_data){
      unsigned idx = (blockDim.x * blockIdx.x) + threadIdx.x;
      if (idx < len_data){
        unsigned sp_a;
        int val = my_data[idx];
        bsearch_range(a, val, len_a, &sp_a);
        my_idx[idx] = sp_a;
        }
    }
    
    
    int main(){
    
    // generate DSIZE random 32-bit integers, PCT_ZERO% are zero
      thrust::host_vector<int> h_data(DSIZE);
      thrust::generate(h_data.begin(), h_data.end(), rand);
      for (int i = 0; i < DSIZE; i++)
        if ((rand()%100)< PCT_ZERO) h_data[i] = 0;
        else h_data[i] %= RNG;
      thrust::device_vector<int> d_data = h_data;
      thrust::device_vector<int> d_result = d_data;
      thrust::sort(d_result.begin(), d_result.end());
      thrust::device_vector<int> d_unique = d_result;
      int unique_size = thrust::unique(d_unique.begin(), d_unique.end()) - d_unique.begin();
      find_my_idx<<< (DSIZE+nTPB-1)/nTPB , nTPB >>>(thrust::raw_pointer_cast(d_unique.data()), unique_size, thrust::raw_pointer_cast(d_data.data()), thrust::raw_pointer_cast(d_result.data()), DSIZE);
    
      thrust::copy(d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
      thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
      return 0;
    }
    $ nvcc t635.cu -o t635
    $ ./t635
    0,6,7,0,3,0,6,0,9,0,0,0,0,9,3,6,0,6,0,6,
    0,2,3,0,1,0,2,0,4,0,0,0,0,4,1,2,0,2,0,2,
    $
    

    【讨论】:

    • 这是一个很好的答案,也是一篇出色的文章。错在我。我遗漏了关于 CCL 对我来说很明显的一点,但这使得这个解决方案不起作用(至少没有一些修改)。查看更新。
    • 我无法轻易找到:Thrust 是否使用 SHUFFLE 进行缩减和前缀扫描?
    • 我不知道thrust是否使用shuffle操作来减少和前缀扫描。 (Thrust 是一个开源库,一个简单的 grep 可能会回答这个问题。)当然,它们不能用于计算能力低于 3.0 的设备。我看不出这个问题与你所描述的问题有什么关系。关于原始问题,随着您的更新,我能想到的解决方案技术涉及直方图作为步骤之一。 (1.) 32 位整数的占用范围或 (2.) 单个 bin 中可能重复的数量是否有任何限制?
    • 我想我更新的答案仍然不太正确,因为输出数据不是“单调递增”。鉴于重复映射到相同值的要求,不确定如何实现。您的问题陈述可能会通过示例输入和所需的输出数据集得到增强,而不是任何额外的措辞。
    • 罗伯特,我真的没想到完整的答案。那会很懒惰,我想要的只是一些提示。我现在知道如何将您的第一个解决方案应用于我的问题(但我需要完整的 3 次通过)。但是您的时间如此慷慨,我认为完全重写该问题会对其他人有所帮助。至于我,我想我知道如何使用前缀扫描和 SHUFL 来做到这一点。我不需要推力。感谢您的帮助。
    【解决方案2】:

    我的答案类似于@RobertCrovella 给出的答案,但我认为使用thrust::lower_bound 而不是自定义二进制搜索更简单。 (既然是纯推力,后端可以互换)

    1. 复制输入数据
    2. 对复制的数据进行排序
    3. 根据已排序的数据创建唯一列表
    4. 在唯一列表中找到每个输入的下限

    我在下面提供了一个完整的示例。有趣的是,通过预先设置排序步骤并再次调用thrust::unique,该过程可以变得更快。根据输入数据,这可以显着减少排序中的元素数量,这是这里的瓶颈。

    #include <iostream>
    #include <stdlib.h>
    #include <thrust/device_vector.h>
    #include <thrust/host_vector.h>
    #include <thrust/transform.h>
    #include <thrust/generate.h>
    #include <thrust/sort.h>
    #include <thrust/unique.h>
    #include <thrust/binary_search.h>
    #include <thrust/copy.h>
    
    int main()
    {
      const int ndata = 20;
      // Generate host input data
      thrust::host_vector<int> h_data(ndata);
      thrust::generate(h_data.begin(), h_data.end(), rand);
      for (int i = 0; i < ndata; i++)
      {
        if ((rand() % 100) < 40)
          h_data[i] = 0;
        else
          h_data[i] %= 10;
      }
    
      // Copy data to the device
      thrust::device_vector<int> d_data = h_data;
      // Make a second copy of the data
      thrust::device_vector<int> d_result = d_data;
      // Sort the data copy
      thrust::sort(d_result.begin(), d_result.end());
      // Allocate an array to store unique values
      thrust::device_vector<int> d_unique = d_result;
      {
        // Compress all duplicates
        const auto end = thrust::unique(d_unique.begin(), d_unique.end());
        // Search for all original labels, in this compressed range, and write their
        // indices back as the result
        thrust::lower_bound(
          d_unique.begin(), end, d_data.begin(), d_data.end(), d_result.begin());
      }
    
      thrust::copy(
        d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
      thrust::copy(d_result.begin(),
                   d_result.end(),
                   std::ostream_iterator<int>(std::cout, ","));
      std::cout << std::endl;
      return 0;
    }
    

    【讨论】:

      猜你喜欢
      • 1970-01-01
      • 2017-12-27
      • 2015-12-05
      • 2013-09-29
      • 1970-01-01
      • 1970-01-01
      • 2015-12-30
      • 1970-01-01
      • 2023-01-07
      相关资源
      最近更新 更多