CUDA 推力仿函数 GMEM 访问：ctor 数据复制与 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); }
}

使用 thrust::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); }
}

使用 thrust::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.

上花费多少寄存器来选择其中之一。

CUDA 推力仿函数 GMEM 访问：ctor 数据复制与 ctor dev ptr arg

CUDA Thrust functor GMEM access: ctor data copy vs ctor dev ptr arg

cuda

global-variables

thrust