thrust::device_vector内部的结构

Question

我有一个结构，如

struct Point
{
 int x;
 int y;
float val;
}

我打算使用这个结构来表示稀疏矩阵(我知道CUSPARSE和CUSP，但我只打算使用推力执行一些测试)，并使用推力算法执行操作。

根据我在CUDA编程教程中所学到的，为了更好地进行内存合并，建议使用数组的结构而不是结构的数组。

如果是这样，那么如果我使用上面提到的结构在device_vector中存储非零(以百万为数量级)，该device_vector在处理推力算法时会在GPU中使用未对齐的内存访问吗？

我之所以这样问，是因为我可能需要访问此device_vector中的不规则步长，并通过传递多个函数对象来执行算法操作。

它会像在数组结构上操作的自定义内核一样高效吗？

谢谢。

Answer 1

假设合并内存访问模式，NVIDIA CUDA设备可以高效地访问4、8和16字节结构。为此，CUDA头定义了可以使用的结构int2、int4、float2、float4等。它们被定义为具有有效的对齐方式，因此我建议使用

typedef int2 Point;

当对这些小结构的数组的所有内存访问都是在warp中跨线程的顺序访问时(例如，合并)，并且每个结构元素中的所有数据都由读/写它的线程使用，那么这种类型的AOS访问是非常有效的。事实上，由于不断增加的内存事务，使用这样的向量结构通常会导致比标量数据访问更高的内存吞吐量。

Thrust提供了zip_iterator，以方便和(编码)效率地操作SOA数据，就好像它是AOS数据一样。因此，虽然小结构在直接CUDA C++中是有效的，但在使用推力时，您可以选择为每个结构成员使用单独的device_vector，并在调用transform和其他推力算法之前使用zip_iterator将它们压缩在一起。推力示例代码中包含了一些这样的示例。