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
1 Like
njuffa
July 28, 2021, 8:19pm
2
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
system
Closed
September 28, 2021, 10:20am
5
This topic was automatically closed 60 days after the last reply. New replies are no longer allowed.