NVIDIA / NVIDIA/cutlass

[FEA] CVT F32 -> TF32 PTX for sm80

Open
#2,254 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage feature request inactive-30d inactive-90d
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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.