Float to TF32 conversion

hello i would like to know if there is any instruction to convert tf32 to float, thanks

i tried this

#include <mma.h>
#include <iostream>
#include <stdio.h>

__global__ void Cvt_FT_F() {
    float a = 1.0;
    uint32_t b;

    // Cvt Float - TF32
    asm("cvt.rna.tf32.f32         %0, %1;\n" : "=r"(b) : "f"(a));
    // Cvt TF32 - Float
    asm("cvt.rna.f32.tf32         %0, %1;\n" : "=f"(a) : "r"(b));

    printf("%f \n", a);
}

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

but i get this error

ptxas /tmp/tmpxft_000016eb_00000000-6_Cvt_FT_F.ptx, line 43; error   : Illegal rounding modifier for instruction 'cvt'
ptxas /tmp/tmpxft_000016eb_00000000-6_Cvt_FT_F.ptx, line 43; error   : Unexpected instruction types specified for 'cvt'
ptxas fatal   : Ptx assembly aborted due to errors

thank for help

Are you building for compute capability >= 8.0 (sm_80) with a fairly recent CUDA version?

The PTX documentation only lists cvt.rna.tf32.f32. I don’t see any mention of conversion in the opposite direction, so an operation cvt.rna.f32.tf32 probably doesn’t exist.

float to tf32 is already provided for, using the first instruction. The second instruction cvt.rna.f32.tf32 is not valid.

tf32 is already compliant with float from a format perspective. You can reinterpret it as a float. Here is an example:

__device__ float cvt_demo(float a){
    float ret;
    asm ("{.reg .b32 %mr;\n"
        "cvt.rna.tf32.f32 %mr, %1;\n"
        "mov.b32 %0, %mr;}\n" : "=f"(ret) : "f"(a));
    return ret;
}
__global__ void Cvt_FT_F(float *b) {
    *b = cvt_demo(*b);
}

int main() {
    float *a;
    cudaMallocManaged(&a, sizeof(float));
    Cvt_FT_F<<<1, 1>>>(a);
    cudaDeviceSynchronize();
}

and, yes, if you compile for a GPU architecture less than cc8.0, you will get compile errors

1 Like

Thank you for your help

This topic was automatically closed 60 days after the last reply. New replies are no longer allowed.