NVIDIA / NVIDIA/cutlass

[BUG] CUTLASS Python Interface nvrtc fails on Hopper

Open
#2,150 2 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

  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 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.