CUDA 推力函子 GMEM 访问:ctor data copy vs ctor dev ptr arg

Question

我有两种方法可以让我的推力函子访问全局不可矢量化非均匀访问的只读状态。不幸的是，内核执行时间相差 100 倍。为什么我的两种策略会有什么不同？

更一般地说：有没有一种规范的方式来为推力函子提供对这些类型的全局变量的访问？

我的

第一种方法是将我的全局数据的副本放入函子中。推力机械似乎在设备上执行上传和缓存：

// functor containing a copy of array dependency
template<size_t BARSIZE>
struct foo1_func
{
  __align__(16) float bar[BARSIZE];
  foo1_func(float _bar[BARSIZE]) { memcpy(bar,_bar,BARSIZE*sizeof(float)); }
  __host__ __device__ operator()(float &t) { t = do_something(t, bar); }
}

调用使用推力：：for_each...

// assuming barData is a float[]
foo<N>(barData);

我的第二种方法是使用 thrust：：copy 自己执行上传到设备，然后将上传数据的设备内存指针传递给我的函子。此方法似乎要慢得多：

// functor containing device pointers to array in GMEM
struct foo2_func
{
  float *bar;
  foo2_func(float* _bar) { bar = bar; }
  __host__ __device__ operator()(float &t) { t = do_something(t, bar); }
}

调用使用推力：：for_each...

// assuming d_bar is a thrust::device_vector
foo(thrust::raw_pointer_cast(d_bar.data()));

指向说明规范或唯一函子模式的源的链接，我们感激地接受了。

Answer 1

使用第一种方法，您实际上是尝试通过将结构foo1_func作为内核函数参数传递来将整个数组bar放入 GPU 寄存器。

__global__ void kernel_generated_by_thrust(struct foo_func f, ...) {
  float x = f.bar[3];
  ...
}

如果bar的尺寸小到可以放入电阻器中，则随机访问bar实际上是随机访问寄存器。

但是您的第二种方法仅通过结构传递全局内存指针。因此，对bar的随机访问是对全局内存的随机访问。

这就是为什么第二种方式要慢得多的原因。

这两种方式都有其用例。您可以根据要实现的目标，bar的大小以及要花费多少寄存器来缓存bar来选择其中任何一个。