[BUG] CUTLASS Python Interface nvrtc fails on Hopper
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Describe the bug
After switching the device compiler in the CUTLASS Python Interface backend to nvrtc, running even simple examples fails.
It seems that nvrtc is missing C++ std type traits used in the CUTLASS 3.x API code generated for GEMM on Hopper.
Making these traits available via the definitions in https://github.com/NVIDIA/jitify allows execution to complete successfully.
Steps/Code to reproduce bug
Optionally, create and activate a virtual environment.
Then, clone this repository and install the CUTLASS Python interface by running the following command in the root directory of the cloned repository:
pip install -e .
Store the following code in a file called example.py:
import cutlass
import numpy as np
cutlass.backend.compiler.nvrtc()
plan = cutlass.op.Gemm(element=np.float16, layout=cutlass.LayoutType.RowMajor)
A, B, C, D = [np.ones((128, 128), dtype=np.float16) for i in range(4)]
plan.run(A, B, C, D)
This is identical to the basic usage example given in python/README.md, with one difference: It switches the Python interface backend device compiler from nvcc to nvrtc via the provided API.
Executing the example with python example.py results in the following compilation error being reported:
RuntimeError: NVRTC Error: 6
cutlass/python/cutlass_library/../../include/cutlass/epilogue/collective/sm90_epilogue_array_tma_warpspecialized.hpp(497): error: namespace "std" has no member "is_pointer_v"
__CUTE_REQUIRES(std::is_pointer_v<TensorMapC>)
^
cutlass/python/cutlass_library/../../include/cutlass/epilogue/collective/sm90_epilogue_array_tma_warpspecialized.hpp(497): error: type name is not allowed
__CUTE_REQUIRES(std::is_pointer_v<TensorMapC>)
^
cutlass/python/cutlass_library/../../include/cutlass/epilogue/collective/sm90_epilogue_array_tma_warpspecialized.hpp(497): error: expected an expression
__CUTE_REQUIRES(std::is_pointer_v<TensorMapC>)
^
cutlass/python/cutlass_library/../../include/cutlass/epilogue/thread/linear_combination_bias_elementwise.h(92): error: namespace "std" has no member "void_t"
struct ElementwiseOpDispatcher<T, std::void_t<typename T::Arguments>> {
^
cutlass/python/cutlass_library/../../include/cutlass/epilogue/thread/linear_combination_bias_elementwise.h(92): error: expected a ">"
struct ElementwiseOpDispatcher<T, std::void_t<typename T::Arguments>> {
^
cutlass/python/cutlass_library/../../include/cutlass/epilogue/thread/linear_combination_bias_elementwise.h(92): error: expected a ";"
struct ElementwiseOpDispatcher<T, std::void_t<typename T::Arguments>> {
^
cutlass/python/cutlass_library/../../include/cutlass/epilogue/thread/linear_combination_bias_elementwise.h(106): warning #12-D: parsing restarts here after previous syntax error
};
^
Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"
6 errors detected in the compilation of "module.cu".
The root cause seems to be missing type traits:
error: namespace "std" has no member "is_pointer_v"
Expected behavior
This example should run without error.
Environment details (please complete the following information):
- OS: Linux, Ubuntu 22.04, kernel 5.15
- Device: Nvidia H100 PCIe
- CUDA version: 12.8
- Nvidia driver version: 550.90.07
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 headers named in the NVRTC errors, especially sm90_epilogue_array_tma_warpspecialized.hpp and linear_combination_bias_elementwise.h, and compare the available definitions with NVIDIA/jitify. Reproduce the failure using the example.py script and the setup steps in python/README.md. Done means the basic GEMM example runs through the nvrtc backend without the missing std::is_pointer_v or std::void_t errors.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp, python
- Domain
- api, compilers
- Issue type
- Bug
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 45/100