[FEA] CVT F32 -> TF32 PTX for sm80
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Is your feature request related to a problem? Please describe.
Currently, converting from tf32 to f32 with round to nearest dispatches to a PTX cvt instruction only for sm90.
Describe the solution you'd like
If we allow rna rounding, we can dispatch to cvt.rna.tf32.f32, which works for sm80.
Describe alternatives you've considered
N/A
Additional context
A simple code sample is given below:
__global__ void f2tfK() {
constexpr float x = -0.45466f;
uint32_t d = 0;
constexpr auto f2tf = cutlass::NumericConverter<cutlass::tfloat32_t, float>{};
asm volatile("cvt.rna.tf32.f32 %0, %1;" : "=r"(d) : "f"(x));
const auto res = cutlass::tfloat32_t::bitcast(d);
const auto cRes = f2tf(x);
printf("Intrinsic: "); cute::print(res); printf("\n");
printf("Other: "); cute::print(cRes); printf("\n");
printf("isEqual? %s\n", cRes == res ? "yes" : "no");
}
// Output:
// Intrinsic: -0.454590
// Other: -0.454590
// isEqual? yes
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start in include/cutlass/numeric_conversion.h around the linked conversion dispatch. Compare the existing sm90 PTX path with the provided cvt.rna.tf32.f32 CUDA sample and determine the sm80 condition for allowing RNA rounding. Done means sm80 conversion matches the intrinsic behavior shown in the sample.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc, performance
- Issue type
- Feature
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100