use*_*757 1 c++ memory precision cuda
half我对定义设备上的精确数据分配的正确方法感到困惑。例如，如果我希望有一个half设备内存数组，是否应该这样分配？
__device__ half array[32];
或者这会与数组的实际设备内存使用量相同，float因为每个half数组之前都存储有未使用的位？（就像 abool必须占用整个内存地址，而不仅仅是 1 位）以下分配的正确方法是吗？
__device__ half2 array[16];
如果两者分配的字节数相同，那么有什么意义呢half2？
两个声明占用相同的空间。从代码正确性的角度来看，这两种方法都应该可用。
half2应该首选至少有两个原因：
当进行扭曲范围的加载时，half2每个线程一个比每个half线程一个更有效的加载（或存储）。
加法和乘法等某些操作只有在half2类型上完成时（在本次讨论中忽略 Volta tensorCore 操作），才能在支持此类操作的架构（cc5.3 和 cc6.1 及更高版本）上实现完整（即最高）吞吐量。
在其他方面，关于是否使用其中一种的决定可能相当于询问任何其他向量类型，例如intvs. int2。我不会尝试在这里对使用向量类型的动机进行完整的总结。
给定的两个示例的底层内存存储模式是相同的，因此一般来说，正确对齐的指针 tohalf应该可以强制转换（例如，在正确/交替索引边界上），half2同样，half2指针应该可以安全地强制转换为half。
即使您选择使用half2，在某些情况下使用 也half可能是适当的或不可避免的，例如，如果您需要修改单个数量，或者例如对于某些 CUBLAS 函数调用。在这种情况下，至少对于最新版本的 CUDA，可以像任何其他向量类型一样转换half2为half（反之亦然）：
__device__ half2 array[16];
...
half2 myval = array[0];
half first = myval.x;
half second = myval.y;
正如已经提到的，half2指针可以安全地转换为half：
half2 *data = array+2;
half *hdata = reinterpret_cast<half *>(data);
这个有点相关的问题可能也很有趣。