如何正确分配cuda半精度阵列?

两个声明占用相同的空间。从代码正确性的角度来看,这两种方法都应该是可用的。

half2 出于至少两个原因,应优先考虑:

1. 进行经纱宽度加载时 半2 每个线程的加载(或存储)效率将高于一个线程 half 每个线程。

2. 一些操作如下 addition and multiplication 只有在 半2 键入(在此讨论中忽略Volta tensorCore操作)。

在其他方面,关于是否使用一个或另一个的决定可能等同于询问任何其他向量类型,例如 int vs。 int2 。我不打算在这里对使用向量类型的动机进行完整的总结。

给定的两个示例的底层内存存储模式是相同的,因此一般来说 一半 应可浇铸(例如,在适当/交替索引边界上),以 半2 同样 半2 指针应该可以安全地投射到 一半 。

即使您选择使用 半2 ,使用 一半 在某些情况下可能是适当的或不可避免的,例如,如果您需要修改单个数量,或例如 certain CUBLAS function calls 。在这种情况下,至少在CUDA的最新版本中,可以转换 半2 到 一半 (反之亦然)与任何其他向量类型一样:

__device__ half2 array[16];
...
half2 myval = array[0];
half first = myval.x;
half second = myval.y;

如前所述 半2 指针可以安全地投射到 一半 :

half2 *data = array+2;
half  *hdata = reinterpret_cast<half *>(data);

This 有些相关的问题可能也很有趣。