NVIDIA / NVIDIA/cutlass

[BUG] Possible Bug when exporting Kernel as PyTorch extension

Open
#1,655 5 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

Description

I have the following script:

`import random
import cutlass
import torch

print_module = False

batch = 256
feature_dim_in = 4098
feature_dim_out = 10

type_A = torch.int8
type_B = torch.int8
type_C = torch.int32
type_D = torch.int32

torch.random.manual_seed(12)
random.seed(22)

input = torch.randint(-3, 3, size=(batch,feature_dim_in), device='cuda').to(type_A)
weight = torch.randint(-3, 3, size=(feature_dim_out,feature_dim_in), device='cuda').to(type_B)
weight_col_major = weight.t().contiguous()#.t()
bias = torch.randint(-3, 3, size=(batch,feature_dim_out), device='cuda').to(type_C)

result_tensor = torch.zeros_like(bias)
#Forward pass:

plan = cutlass.op.Gemm(
alpha=1, beta=1,
element_A=type_A,
element_B=type_B,
element_C=type_C,
element_D=type_D,
layout_A=cutlass.LayoutType.RowMajor,
layout_B=cutlass.LayoutType.RowMajor,
layout_C=cutlass.LayoutType.RowMajor
)
output_torch = (torch.matmul(input.float(),weight.t().float()) + bias.float()).to(torch.int32)
plan.opclass = cutlass.OpcodeClass.Simt
plan.run(input, weight_col_major, bias, result_tensor, print_module=print_module)

op = plan.construct()
linear_layer_gemm = cutlass.emit.pytorch(op, name='linear_layer', cc=plan.cc, sourcedir='linear', jit=True)`

when I call the python script I receive following error Traceback:

`Traceback (most recent call last):
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/utils/cpp_extension.py", line 2107, in _run_ninja_build
subprocess.run(
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/subprocess.py", line 516, in run
raise CalledProcessError(retcode, process.args,
subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1.

The above exception was the direct cause of the following exception:

Traceback (most recent call last):
File "linear_layer.py", line 55, in
linear_layer_gemm = cutlass.emit.pytorch(op, name='linear_layer', cc=plan.cc, sourcedir='linear', jit=True)
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass/emit/pytorch.py", line 927, in pytorch
return _pytorch_gemm(device_op, name, cc, jit, sourcedir)
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass/emit/pytorch.py", line 775, in _pytorch_gemm
return _jit(name, cc, cpp_file, cuda_file)
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass/emit/pytorch.py", line 697, in jit
jitmodule = load(
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/utils/cpp_extension.py", line 1309, in load
return jit_compile(
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/utils/cpp_extension.py", line 1719, in jit_compile
write_ninja_file_and_build_library(
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/utils/cpp_extension.py", line 1832, in write_ninja_file_and_build_library
run_ninja_build(
File "/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/utils/cpp_extension.py", line 2123, in run_ninja_build
raise RuntimeError(message) from e
RuntimeError: Error building extension 'linear_layer': [1/3] c++ -MMD -MF linear_layer.o.d -DTORCH_EXTENSION_NAME=linear_layer -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE="gcc" -DPYBIND11_STDLIB="libstdcpp" -DPYBIND11_BUILD_ABI="cxxabi1011" -I/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass_library/source/include -I/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass_library/source/tools/util/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/torch/csrc/api/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/TH -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/THC -isystem /home/miniconda3/envs/cuda11-8/include -isystem /home/miniconda3/envs/cuda11-8/include/python3.8 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -c /home/reproai/reprai/extensions/linear/linear_layer.cpp -o linear_layer.o
[2/3] /home/miniconda3/envs/cuda11-8/bin/nvcc --generate-dependencies-with-compile --dependency-output linear_layer_kernel.cuda.o.d -DTORCH_EXTENSION_NAME=linear_layer -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE="gcc" -DPYBIND11_STDLIB="libstdcpp" -DPYBIND11_BUILD_ABI="cxxabi1011" -I/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass_library/source/include -I/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass_library/source/tools/util/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/torch/csrc/api/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/TH -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/THC -isystem /home/miniconda3/envs/cuda11-8/include -isystem /home/miniconda3/envs/cuda11-8/include/python3.8 -D_GLIBCXX_USE_CXX11_ABI=0 -D__CUDA_NO_HALF_OPERATORS -D__CUDA_NO_HALF_CONVERSIONS
-D__CUDA_NO_BFLOAT16_CONVERSIONS
-D__CUDA_NO_HALF2_OPERATORS
--expt-relaxed-constexpr -gencode=arch=compute_89,code=sm_89 --compiler-options '-fPIC' -std=c++17 -c /home/reproai/reprai/extensions/linear/linear_layer_kernel.cu -o linear_layer_kernel.cuda.o
FAILED: linear_layer_kernel.cuda.o
/home/miniconda3/envs/cuda11-8/bin/nvcc --generate-dependencies-with-compile --dependency-output linear_layer_kernel.cuda.o.d -DTORCH_EXTENSION_NAME=linear_layer -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE="gcc" -DPYBIND11_STDLIB="libstdcpp" -DPYBIND11_BUILD_ABI="cxxabi1011" -I/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass_library/source/include -I/home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/cutlass_library/source/tools/util/include -isystem /home/miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include -isystem /home//miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/torch/csrc/api/include -isystem /home//miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/TH -isystem /home//miniconda3/envs/cuda11-8/lib/python3.8/site-packages/torch/include/THC -isystem /home//miniconda3/envs/cuda11-8/include -isystem /home//miniconda3/envs/cuda11-8/include/python3.8 -D_GLIBCXX_USE_CXX11_ABI=0 -D__CUDA_NO_HALF_OPERATORS -D__CUDA_NO_HALF_CONVERSIONS
-D__CUDA_NO_BFLOAT16_CONVERSIONS
-D__CUDA_NO_HALF2_OPERATORS
--expt-relaxed-constexpr -gencode=arch=compute_89,code=sm_89 --compiler-options '-fPIC' -std=c++17 -c /home//reproai/reprai/extensions/linear/linear_layer_kernel.cu -o linear_layer_kernel.cuda.o
/home//reproai/reprai/extensions/linear/linear_layer_kernel.cu(106): error: namespace "torch" has no member "I32"

1 error detected in the compilation of "/home//reproai/reprai/extensions/linear/linear_layer_kernel.cu".
ninja: build stopped: subcommand failed.
`

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 cutlass/emit/pytorch.py, especially the _pytorch_gemm and _jit paths named in the traceback, then inspect generated linear_layer_kernel.cu at line 106. Reproduce the supplied script and verify that the exported PyTorch extension compiles without the reported torch.I32 error.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python, pytorch
Domain
build-system, machine-learning
Issue type
Bug
Difficulty
3/5
Estimated time
1-2 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.