NVIDIA / NVIDIA/cutlass

how to use nvrtc to run a sm90 kernel

Open
#1,889 11 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-30d inactive-90d question
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Hi

I want to use nvrtc to compile a sm90 kernel in runtime. The problem is that I don't have the kernel instance on host thus can't run to_underlying_arguments to get kernel param to launch the kernel.

When running a sm80 code, I can get it's params directly since Params is equal to Arguments, but for a sm90 kernel I needs to construct tma_load struct before I launch the kernel.

I see there is a nvrtc example in cutlass which use python to launch a nvcc progress to compile host api for the kernel such as initialize into a .so. But I have a c++ application and can't launch any child progress in it.

So is there anyway to runtime compile a sm90 kernel in my application. I see make_tma_copy is defined as CUDA_HOST_RTC which can be used in nvrtc, but it will call cuTensorMapEncodeTiled which can only called on host, and nvrtcCompileProgram can only compile a kernel function. Is there any example to use make_tma_copy in nvrtc, thanks.

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 with the CUTLASS NVRTC example and trace how it handles host-side setup. Compare the sm80 argument path with run_to_underlying_arguments, make_tma_copy, cuTensorMapEncodeTiled, and nvrtcCompileProgram. Done means a documented C++ runtime-compilation path for sm90, or a clearly stated limitation if the requested flow is unsupported.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python
Domain
hpc, tooling
Issue type
Feature
Difficulty
5/5
Estimated time
Over a week
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
20/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.