【问题标题】:How to avoid default construction of elements in thrust::device_vector?如何避免在thrust::device_vector 中默认构造元素?
【发布时间】:2013-04-29 16:08:04
【问题描述】:
  1. 似乎在创建新的推力向量时所有元素默认为 0 - 我只是想确认这将永远是这种情况。

  2. 如果是这样,是否还有一种方法可以绕过负责此行为的构造函数以提高速度(因为对于某些向量,我不需要它们具有初始值,例如如果它们的原始指针作为输出传递给 CUBLAS)?

【问题讨论】:

    标签: c cuda gpgpu thrust


    【解决方案1】:

    thrust::device_vector 使用其提供的分配器构造它包含的元素,就像std::vector 一样。当向量要求它构造一个元素时,可以控制分配器做什么。

    使用自定义分配器来避免向量元素的默认初始化:

    // uninitialized_allocator is an allocator which
    // derives from device_allocator and which has a
    // no-op construct member function
    template<typename T>
      struct uninitialized_allocator
        : thrust::device_malloc_allocator<T>
    {
      // note that construct is annotated as
      // a __host__ __device__ function
      __host__ __device__
      void construct(T *p)
      {
        // no-op
      }
    };
    
    // to make a device_vector which does not initialize its elements,
    // use uninitialized_allocator as the 2nd template parameter
    typedef thrust::device_vector<float, uninitialized_allocator<float> > uninitialized_vector;
    

    调用uninitialized_allocator::construct 仍会产生内核启动成本,但该内核将是一个空操作,会很快退出。您真正感兴趣的是避免填充数组所需的内存带宽,此解决方案就是这样做的。

    有一个完整的示例代码here

    请注意,此技术需要 Thrust 1.7 或更高版本。

    【讨论】:

    • 非常好。尽管之前为 stl 编写了调试分配器重载,但我忘记了最终的构造调用在这里。应该一直挖。 +1 :)
    • 实际上——我很困惑。我错了,您的示例中的resize 链接到insert,而后者又链接到fill_insert,最后又是uninitialized_fill_n?所以你仍然得到副本,尽管在那里设置新的storage_type 区域时可能忽略了construct? ...显然我需要在调试器中逐步执行此操作,但我没有看到如何避免来自默认/值初始化 x 默认 arg 的最终 uninitialized_fill_n
    • 您可能需要在具有最新推力的调试器中单步执行。这是一个复杂的调度。
    • 知道了,抱歉。出于某种原因,我只是错过了 vector_base.inl 中 resize() 的单参数版本。打扰了!对于任何观看的人来说,堆栈是:(1) detail::contiguous_storage::contiguous_storage(unsigned int n, const uninitialized_allocator & alloc) (2) detail::vector_base::append(unsigned int n) (3) detail::vector_base::resize(unsigned int new_size)
    猜你喜欢
    • 2011-11-05
    • 2012-02-21
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 2011-08-21
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多