how to use nvrtc to run a sm90 kernel
Nobody has claimed this yet.
- 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
- 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 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