半精度:__float2half与__float2half_rn之间的差异

Question

似乎没有关于这两个功能的文档.

__float2half和之间有什么区别__float2half_rn？

Answer 1

看来CUDA文档在这里确实有点不足.

在CUDA 7.5中引入新数据类型之前unsigned short __float2half_rn(float),结合使用的函数float __half2float(unsigned short x)已经存在于CUDA half中.它的定义是device_functions.h.那里的评论如下:

将单精度浮点值x转换为以无符号短格式表示的半精度浮点值,采用舍入到最接近均匀模式.

该函数half __float2half(float)定义在cuda_fp16.h并且显然相同,但返回一个half:

浮点数转换一个半精度舍入到最近的模式.

但是,由于half是一个typedef unsigned short,我检查他们是否也这样做,使用以下代码:

#include <stdio.h>
#include "cuda_fp16.h"
#include "device_functions.h"
__global__ void test()
{
//  auto test = __float2half( 1.4232 );
    auto test = __float2half_rn( 1.4232 );
    printf( "%hu\n", test );
}

int main()
{
    test<<<1,1>>>();
    cudaDeviceSynchronize();
}

我发现(for sm_20)old __float2half_rn()有一个额外的int16到int32操作并且有一个32位存储.另一方面,__float2half_()没有这个转换并且做了16位商店.

相关的SASS代码__float2half_rn():

/*0040*/         I2I.U32.U16 R0, R0;
/*0050*/         STL [R2], R0;

用于__float2half():

/*0048*/         STL.U16 [R2], R0;