cuda - 在CUDA中使用Half2

我正在尝试使用half2，但遇到错误，即

error: class "__half2" has no member "y"

发生错误的代码部分如下：

uint8_t V_ [128];       // some elements (uint8), to save space
float   V_C[128];       // storing the diff to use later
half2 *C_ = C.elements; // D halfs stored as half2, to be read
Cvalue = 0.0;
for (d = 0; d < D; d+=2)
{
  V_C [d  ] = V_[d]   - __half2float(C_[d/2].x)    ;
  V_C [d+1] = V_[d+1] - __half2float(C_[d/2].y)    ;
  Cvalue   += V_C [d]   * V_C [d]  ;
  Cvalue   += V_C [d+1] * V_C [d+1];
}

有什么帮助吗？

更新：
谢谢您的帮助！我终于使用了以下内容...

uint8_t V_ [128] ;
float   V_C[128] ;
const half2 *C_ = C.elements;
Cvalue = 0.0;
float2 temp_;
for (d = 0; d < D; d+=2)
  {
    temp_     = __half22float2(C_[d/2]);
    V_C [d  ] = V_[d]   - temp_.x      ;
    V_C [d+1] = V_[d+1] - temp_.y      ;
    Cvalue   += V_C [d]   * V_C [d]  ;
    Cvalue   += V_C [d+1] * V_C [d+1];
  }

我的特定应用程序略有加速，因为全局内存的负载是瓶颈。

最佳答案

您不能使用点运算符访问half2的某些部分，应该为此使用内部函数。

从documentation：

__CUDA_FP16_DECL__ float __high2float ( const __half2 a )
    Converts high 16 bits of half2 to float and returns the result.
__CUDA_FP16_DECL__ __half __high2half ( const __half2 a )
    Returns high 16 bits of half2 input.
__CUDA_FP16_DECL__ __half2 __high2half2 ( const __half2 a )
    Extracts high 16 bits from half2 input.
__CUDA_FP16_DECL__ __half2 __highs2half2 ( const __half2 a, const __half2 b )
    Extracts high 16 bits from each of the two half2 inputs and combines into one half2 number.
__CUDA_FP16_DECL__ float __low2float ( const __half2 a )
    Converts low 16 bits of half2 to float and returns the result.
__CUDA_FP16_DECL__ __half __low2half ( const __half2 a )
    Returns low 16 bits of half2 input.
__CUDA_FP16_DECL__ __half2 __low2half2 ( const __half2 a )
    Extracts low 16 bits from half2 input.
__CUDA_FP16_DECL__ __half2 __lowhigh2highlow ( const __half2 a )
    Swaps both halves of the half2 input.
__CUDA_FP16_DECL__ __half2 __lows2half2 ( const __half2 a, const __half2 b )
    Extracts low 16 bits from each of the two half2 inputs and combines into one half2 number.

不仅如此，这行取决于C.elements是什么类型

half2 *C_ = C.elements; // D halfs stored as half2, to be read

可能是错误的（如果C.elements是half*。此处的注释不明确）。
half2不是一对half。
实际上，在当前实现中，half2只是包装在结构中的unsigned int：

// cuda_fp16.h

typedef struct __align__(2) {
   unsigned short x;
} __half;

typedef struct __align__(4) {
   unsigned int x;
} __half2;

#ifndef CUDA_NO_HALF
typedef __half half;
typedef __half2 half2;
#endif /*CUDA_NO_HALF*/

没有人说可以将half数组作为half2数组访问。

关于cuda - 在CUDA中使用Half2，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/37133128/