在CUDA中使用half2

Question

我试图使用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];
  }

我在我的特定应用程序中获得了轻微的加速,因为来自全局内存的负载是瓶颈......

Answer 1

你不能访问half2带点运算符的部分,你应该使用内部函数.

从文档:

__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是a half*.这里的评论不清楚). half2不是一对halfs.实际上,在当前的实现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可以将s 数组作为s 数组访问half2.