时间戳

首页 > 解决方案 > CUDA 会自动将 float4 数组转换为数组结构吗？

cuda - CUDA 会自动将 float4 数组转换为数组结构吗？

问题描述

我有以下代码片段：

#include <stdio.h>

struct Nonsense {
    float3 group;
    float other;
};

__global__ void coalesced(float4* float4Array, Nonsense* nonsenseArray) {
    float4 someCoordinate = float4Array[threadIdx.x];
    someCoordinate.x = 5;
    float4Array[threadIdx.x] = someCoordinate;

    Nonsense nonsenseValue = nonsenseArray[threadIdx.x];
    nonsenseValue.other = 3;
    nonsenseArray[threadIdx.x] = nonsenseValue;
}

int main() {
    float4* float4Array;
    cudaMalloc(&float4Array, 32 * sizeof(float4));
    cudaMemset(float4Array, 32 * sizeof(float4), 0);

    Nonsense* nonsenseArray;
    cudaMalloc(&nonsenseArray, 32 * sizeof(Nonsense));
    cudaMemset(nonsenseArray, 32 * sizeof(Nonsense), 0);

    coalesced<<<1, 32>>>(float4Array, nonsenseArray);
    cudaDeviceSynchronize();
    return 0;
}

当我通过 Nsight 中的 Nvidia 分析器运行它并查看全局内存访问模式时，float4Array 具有完美的合并读取和写入。同时，Nonsense 数组的访问模式很差（因为它是一个结构数组）。

NVCC 是否会自动将 float4 数组（在概念上是结构数组）转换为数组结构，以获得更好的内存访问模式？

标签： cudansight

解决方案

不，它不会将其转换为数组结构。我认为，如果您仔细考虑这一点，您会得出结论，编译器几乎不可能以这种方式重新组织数据。毕竟，被传递的东西是一个指针。

只有一个数组，并且该数组的元素仍然具有相同顺序的结构元素：

float address (i.e. index):      0      1      2      3      4      5 ...
array element             : a[0].x a[0].y a[0].z a[0].w a[1].x a[1].y ...

然而，该float4数组提供了更好的模式，因为编译器会为每个线程生成一个 16 字节的加载。这有时被称为“向量加载”，因为我们正在为float4每个线程加载一个向量（在这种情况下）。因此，相邻线程仍在读取相邻数据，并且您具有理想的合并行为。在上面的例子中，线程 0 会读取a[0].x,和a[0].y,线程 1 会读取等等。所有这些都将发生在单个请求（即 SASS 指令）中，但可能会被拆分到多个事务中。将请求拆分为多个事务不会导致任何效率损失（在这种情况下）。a[0].za[0].wa[1].xa[1].y

在Nonsense结构的情况下，编译器无法识别该结构也可以以类似的方式加载，因此在后台它必须为每个线程生成 3 或 4 次加载：

  • 一个 8 字节加载（或两个 4 字节加载）加载的前两个字float3 group
  • 一个 4 字节加载加载的最后一个字float3 group
  • 一个 4 字节的负载来加载float other

如果您绘制每个线程的上述负载，也许使用上图，您会看到每个负载都涉及一个步幅（每个线程加载的项目之间的未使用元素），因此会导致效率降低。

通过在结构中使用仔细的类型转换或联合定义，您可以让编译器Nonsense在一次加载中加载您的结构。

这个答案还涵盖了一些与 AoS -> SoA 转换相关的想法以及相关的效率提升。

此答案涵盖矢量负载详细信息。

推荐阅读