__half22float2

// 使用小向量可提升数据搬运效率
__simt_vf__ __launch_bounds__(1024) inline void simt_half22float2(__gm__ half2* input, __gm__ float2* output, uint32_t input_total_length)
{
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 每个线程处理1个half2类型的数据，即2个half类型的数据，因此idx >= input_total_length / 2的线程不处理数据
    if (idx >= input_total_length / 2) {
        return;
    } 
    output[idx] = __half22float2(input[idx]);
}

__global__ __vector__ void cast_kernel(__gm__ half* input, __gm__ float* output, uint32_t input_total_length)
{
    asc_vf_call<simt_half22float2>(dim3(1024), (__gm__ half2*)input, (__gm__ float2*)output, input_total_length);
}