Skip to content

Commit eb56ba2

Browse files
kainzhongpre-commit-ci[bot]janekb04
authored
[Common] Experimental CuTeDSL MXFP8 backends in C++ via TVM-FFI (#3137)
* start draft Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove benchmark scripts Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * add license Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make cutlass dsl required Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix linting errors Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * warn failed compilation Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make tvm-ffi common dependency Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * use a higher version for cutlass-dsl Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * less comment Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * skip zeroing buffer if noop flag set Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * refactor activations Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * lint Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * maybe it's better to prepare tvm-ffi & cutedsl kernels in common/__init__.py Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * also flush colwise scale in SMEM Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * use explicit fp8 types Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * add checks Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * refactor Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fi Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * skip comparision test if cutedsl disabled Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * Merge pull request #3 from janekb04/cutedsl_mxfp8_common Perform noop tensor check on device Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * add test to the qa script Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * warn only once if dlopen fails Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * let tests cover swizzling Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * warn not chosen cutedsl kernel during testing Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * improve lazyload so now the cache is uint32 Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * let tests be verbose about warnings so we know if cutedsl kernel is not dispatched Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * support swizzled scale for specialized kernel but disable it Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix inline ptx Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * more fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * revert to mul Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * add jax tests Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * also ignore noop when dbias Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * relax dim check for the kernels Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * reject 2D quant Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * run more tests with CuTeDSL path Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * make rowwise only handle divisible cases Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fallback to CUDA if not TMA aligned Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * specialize skip masking version Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * add type annotations Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * optimize masking Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * skip masking for some acts Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * allow user to opt out from building with CuTeDSL Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * better warning Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * doc Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * add jax test coverage for cutedsl Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix device sm query Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * catch more python error Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make tvm-ffi required for python users Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * optimize dbias Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * skip scale bound check whenever possible Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * allow colwise to fuse relu too Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * dispatch to rowwise specialized with swizzled scale Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * don't link tvmffi if building a standalong C++ lib without framework Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * change some logging Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * use the right fast math flag Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * also let test script use NVTE_DEBUG Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * also let mutex cover the loading of tvm-ffi func Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * locate tvm-ffi include dir from setup_common_extension Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix how tests run with cutedsl enabled Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * reduce framework scope tests run with cutedsl backend Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * allow C++ tests to run with cutdsl Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove tvmffi and cutlass dependencies from setup.py due to being optional Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * encode sm numbers in tvm_ffi function name & cache key Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make C++ tests run in sequential Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make CuTeDSL default and cover C++ in L1 Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * Improvements and fix to CuTe DSL kernel launches (#7) * [Common] Bypass Python wrapper overhead with __tvm_ffi_object__ Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Avoid heap allocations in DLTensorWrapper Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Avoid querying device in DLTensorWrapper Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Fix packed dtype handling in DLTensorWrapper Signed-off-by: Jan Bielak <jbielak@nvidia.com> * Apply review suggestions Signed-off-by: Jan Bielak <jbielak@nvidia.com> --------- Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * fix dispatching so it never uses specialized if grid won't fit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * also add tests Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * Minor fixes to CuTe DSL infrastructure (#8) * Improvements and fix to CuTe DSL kernel launches (#7) * [Common] Bypass Python wrapper overhead with __tvm_ffi_object__ Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Avoid heap allocations in DLTensorWrapper Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Avoid querying device in DLTensorWrapper Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Fix packed dtype handling in DLTensorWrapper Signed-off-by: Jan Bielak <jbielak@nvidia.com> * Apply review suggestions Signed-off-by: Jan Bielak <jbielak@nvidia.com> --------- Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * [Common] Add missing dtypes to CuTe DSL dispatch Signed-off-by: Jan Bielak <jbielak@nvidia.com> * [Common] Add line information to CuTe DSL inline PTX helpers Signed-off-by: Jan Bielak <jbielak@nvidia.com> --------- Signed-off-by: Jan Bielak <jbielak@nvidia.com> Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * add missing ip & loc Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * refactor so we don't need #pragma to avoid collision in tests Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * whoops Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * whoops again Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * rename NVTE_ENABLE_CUTEDSL_QUANT_BACKEND to NVTE_ENABLE_CUTEDSL_BACKEND Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * switch back to NVTE_WARN_IF_CUTEDSL_BACKEND_NOT_CHOSEN Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * use TE's to_string to convert dtype to str Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove MXFP8QuantFused Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * move python init from tests to TE lib Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * software path for e8m0 conversion Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove unreachable branch Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fold quantize_mxfp8_cutedsl_config.h back to quantize_mxfp8_cutedsl.h Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit for some C++ Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * rewrite dsl_user_op to cute.jit if possible and let compiler optimize Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix python init Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * prefere inline_ptx instead of llvm.inline_asm Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix activation encoding Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove unused device_is_blackwell function Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * link C++ test executable to libtvm_ffi.so Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * move cutedsl utils from math.h to a separate file Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * recover test_cast_mxfp8.cu Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * no need to init extension now Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * hoist importorskip in cutedsl python tests Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix activation_func_to_enum so it works with CUDA 12.8 Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make NVTE_WITH_CUTEDSL both a env var and a cmake option Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix TE path searching for python bootstrap Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * when enabling cutedsl manually, also initialize python if it's not done before Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * use dlsym Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * let main TE lib link to Python::Module Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * add a test to cover untested manually enable cutedsl path Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * update min ver workflow Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * rename L1 CUDA C++ UT Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix build issues for CuTeDSL backend (#10) * disable CuTeDSL support and other fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove docs that mention CuTeDSL since they are off by default now Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * remove test_cutedsl_backend.cpp Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * make envvar.rst more human Signed-off-by: Kaining Zhong <kainingz@nvidia.com> --------- Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix python tests Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * skip L1 C++ UT with CUDA when CuTeDSL is not enabled globally Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix C++ UT python discovery Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix deadlock for multigpu Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix python init Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * fix python init again thanks to codex Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * refactor the serialized initial launch Signed-off-by: Kaining Zhong <kainingz@nvidia.com> * nit Signed-off-by: Kaining Zhong <kainingz@nvidia.com> --------- Signed-off-by: Kaining Zhong <kainingz@nvidia.com> Signed-off-by: Jan Bielak <jbielak@nvidia.com> Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> Co-authored-by: Jan Bielak <janekb04@icloud.com> Co-authored-by: Jan Bielak <jbielak@nvidia.com>
1 parent 14f85fa commit eb56ba2

24 files changed

Lines changed: 5396 additions & 17 deletions

File tree

‎.github/workflows/minimum_versions.yml‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ jobs:
2727
apt-get install -y --allow-change-held-packages git python3.10 python3-pip cudnn9-cuda-12 libnccl-dev libnccl2
2828
python3.10 --version
2929
python3.10 -m pip install torch==${MIN_TORCH_VERSION}
30-
python3.10 -m pip install cmake pybind11[global] ninja pydantic "importlib-metadata>=1.0" packaging einops onnxscript "nvidia-cudnn-frontend>=1.25.0"
30+
python3.10 -m pip install cmake pybind11[global] ninja pydantic "importlib-metadata>=1.0" packaging einops onnxscript "nvidia-cudnn-frontend>=1.25.0" "apache-tvm-ffi>=0.1.12"
3131
- name: 'Checkout'
3232
uses: actions/checkout@v3
3333
with:
@@ -65,7 +65,7 @@ jobs:
6565
apt-get install -y git python3.10 python3-pip cudnn9-cuda-12
6666
python3.10 --version
6767
python3.10 -m pip install jax==${MIN_JAX_VERSION} "flax>=0.7.1"
68-
python3.10 -m pip install cmake pybind11[global] ninja packaging pydantic "nvidia-cudnn-frontend>=1.25.0"
68+
python3.10 -m pip install cmake pybind11[global] ninja packaging pydantic "nvidia-cudnn-frontend>=1.25.0" "apache-tvm-ffi>=0.1.12"
6969
- name: 'Checkout'
7070
uses: actions/checkout@v3
7171
with:

‎build_tools/wheel_utils/build_wheels.sh‎

Lines changed: 20 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@ BUILD_JAX=${5:-true}
1212
CUDA_MAJOR=${6:-12}
1313

1414
export NVTE_RELEASE_BUILD=1
15+
export NVTE_WITH_CUTEDSL=${NVTE_WITH_CUTEDSL:-0}
1516
export PIP_CONSTRAINT=""
1617
export TARGET_BRANCH=${TARGET_BRANCH:-}
1718
mkdir -p /wheelhouse/logs
@@ -23,7 +24,7 @@ git checkout $TARGET_BRANCH
2324
git submodule update --init --recursive
2425

2526
# Install deps
26-
/opt/python/cp310-cp310/bin/pip install cmake pybind11[global] ninja setuptools wheel 'nvidia-cudnn-frontend>=1.25.0'
27+
/opt/python/cp310-cp310/bin/pip install cmake pybind11[global] ninja setuptools wheel 'nvidia-cudnn-frontend>=1.25.0' 'apache-tvm-ffi>=0.1.12'
2728

2829
if $BUILD_METAPACKAGE ; then
2930
cd /TransformerEngine
@@ -35,23 +36,36 @@ if $BUILD_COMMON ; then
3536
VERSION=`cat build_tools/VERSION.txt`
3637
WHL_BASE="transformer_engine-${VERSION}"
3738

39+
if [[ "$NVTE_WITH_CUTEDSL" != "0" ]]; then
40+
WHEEL_PYTHON_TAG=cp310
41+
WHEEL_ABI_TAG=abi3
42+
else
43+
WHEEL_PYTHON_TAG=py3
44+
WHEEL_ABI_TAG=none
45+
fi
46+
3847
# Create the wheel.
39-
/opt/python/cp310-cp310/bin/python setup.py bdist_wheel --verbose --python-tag=py3 --plat-name=$PLATFORM 2>&1 | tee /wheelhouse/logs/common.txt
48+
/opt/python/cp310-cp310/bin/python setup.py bdist_wheel --verbose --python-tag=$WHEEL_PYTHON_TAG --plat-name=$PLATFORM 2>&1 | tee /wheelhouse/logs/common.txt
4049

4150
# Repack the wheel for specific cuda version.
4251
/opt/python/cp310-cp310/bin/wheel unpack dist/*
4352
# From python 3.10 to 3.11, the package name delimiter in metadata got changed from - (hyphen) to _ (underscore).
4453
sed -i "s/Name: transformer-engine/Name: transformer-engine-cu${CUDA_MAJOR}/g" "transformer_engine-${VERSION}/transformer_engine-${VERSION}.dist-info/METADATA"
4554
sed -i "s/Name: transformer_engine/Name: transformer_engine_cu${CUDA_MAJOR}/g" "transformer_engine-${VERSION}/transformer_engine-${VERSION}.dist-info/METADATA"
4655
mv "${WHL_BASE}/${WHL_BASE}.dist-info" "${WHL_BASE}/transformer_engine_cu${CUDA_MAJOR}-${VERSION}.dist-info"
47-
# Set WHEEL Tag to match the py3-none filename written by line 54.
48-
sed -i "s/Tag: cp310-cp310/Tag: py3-none/g" "${WHL_BASE}/transformer_engine_cu${CUDA_MAJOR}-${VERSION}.dist-info/WHEEL"
56+
if [[ "$NVTE_WITH_CUTEDSL" != "0" ]]; then
57+
# Build with CuTeDSL support: uses the CPython 3.10+ Stable ABI.
58+
sed -i "s/Tag: cp310-cp310/Tag: cp310-abi3/g" "${WHL_BASE}/transformer_engine_cu${CUDA_MAJOR}-${VERSION}.dist-info/WHEEL"
59+
else
60+
# Build without CuTeDSL support: remain python version agnostic.
61+
sed -i "s/Tag: cp310-cp310/Tag: py3-none/g" "${WHL_BASE}/transformer_engine_cu${CUDA_MAJOR}-${VERSION}.dist-info/WHEEL"
62+
fi
4963
/opt/python/cp310-cp310/bin/wheel pack ${WHL_BASE}
5064

51-
# Rename the wheel to make it python version agnostic.
65+
# Rename the wheel to match the tag written above.
5266
whl_name=$(basename dist/*)
5367
IFS='-' read -ra whl_parts <<< "$whl_name"
54-
whl_name_target="${whl_parts[0]}_cu${CUDA_MAJOR}-${whl_parts[1]}-py3-none-${whl_parts[4]}"
68+
whl_name_target="${whl_parts[0]}_cu${CUDA_MAJOR}-${whl_parts[1]}-${WHEEL_PYTHON_TAG}-${WHEEL_ABI_TAG}-${whl_parts[4]}"
5569
rm -rf $WHL_BASE dist
5670
mv *.whl /wheelhouse/"$whl_name_target"
5771
fi

‎docs/envvars.rst‎

Lines changed: 24 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -93,6 +93,15 @@ Optional Dependencies
9393
:Default: ``0``
9494
:Description: Enable NVSHMEM support. When set to ``1``, requires ``NVSHMEM_HOME`` to be set to the NVSHMEM installation directory.
9595

96+
.. envvar:: NVTE_WITH_CUTEDSL
97+
98+
:Type: ``int`` (0 or 1)
99+
:Default: ``0``
100+
:Description: Enable CuTeDSL support when building ``libtransformer_engine.so``. When set to ``1``, requires
101+
CPython Stable ABIs from Python 3.10 and above, and ``apache-tvm-ffi`` installed.
102+
Note: when linking ``libtransformer_engine.so`` built with ``NVTE_WITH_CUTEDSL=1`` to a standalone C++ executable,
103+
you must provide the definition of these CPython symbols. See ``tests/cpp/CMakeLists.txt`` for an example.
104+
96105
.. envvar:: NVTE_BUILD_ACTIVATION_WITH_FAST_MATH
97106

98107
:Type: CMake option
@@ -136,6 +145,20 @@ General
136145
value if an application legitimately creates more tensor handles than
137146
the default pool can hold.
138147

148+
.. envvar:: NVTE_ENABLE_CUTEDSL_BACKEND
149+
150+
:Type: ``int`` (0 or 1)
151+
:Default: ``0``
152+
:Description: Use CuTeDSL kernel implementations when available. This requires that Transformer Engine was built with ``NVTE_WITH_CUTEDSL=1``.
153+
It also requires ``apache-tvm-ffi``, ``nvidia-cutlass-dsl`` to be installed and you are using CPython 3.10 or later.
154+
If these requirements are not met, Transformer Engine will fall back to using CUDA C++ kernels if possible.
155+
156+
.. envvar:: NVTE_WARN_IF_CUTEDSL_BACKEND_NOT_CHOSEN
157+
158+
:Type: ``int`` (0 or 1)
159+
:Default: ``0``
160+
:Description: Warn when TE falls back to the CUDA C++ kernels instead of dispatching to available CuTeDSL kernels. Useful to verify if the CuTeDSL path is taken.
161+
139162
.. envvar:: NVTE_GROUPED_TENSOR_HANDLE_POOL_SIZE_MB
140163

141164
:Type: ``int`` (positive integer)
@@ -311,7 +334,7 @@ Kernel Configuration
311334

312335
:Type: ``int`` (0 or 1)
313336
:Default: ``0``
314-
:Description: Enable fast math optimizations in runtime-compiled (NVRTC) kernels. This trades numerical accuracy for performance. These optimizations are experimental and inconsistently implemented.
337+
:Description: Enable fast math optimizations in runtime-compiled (NVRTC / CuTeDSL) kernels. This trades numerical accuracy for performance. These optimizations are experimental and inconsistently implemented.
315338

316339
.. envvar:: NVTE_DISABLE_NVRTC
317340

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
# Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2+
#
3+
# See LICENSE for license information.
4+
5+
set -e
6+
7+
: ${XML_LOG_DIR:=/logs}
8+
mkdir -p "$XML_LOG_DIR"
9+
10+
# Find TE
11+
: ${TE_PATH:=/opt/transformerengine}
12+
TE_LIB_PATH=$(pip3 show transformer-engine | grep -E "Location:|Editable project location:" | tail -n 1 | awk '{print $NF}')
13+
export LD_LIBRARY_PATH=$TE_LIB_PATH:$LD_LIBRARY_PATH
14+
15+
# Set parallelization parameters
16+
NUM_PHYSICAL_CORES=$(nproc)
17+
NUM_PARALLEL_JOBS=4
18+
19+
# If NVTE_ENABLE_CUTEDSL_BACKEND is not set to 1, we assume this also applies to qa/L0_cppunittest,
20+
# in which case we don't need to run this test since L0_cppunittest already covers the CUDA path.
21+
# Otherwise, we override NVTE_ENABLE_CUTEDSL_BACKEND to 0 to force the test to take the CUDA implementation path,
22+
# so we can cover both CUDA and CuteDSL paths on CI.
23+
if [ "${NVTE_ENABLE_CUTEDSL_BACKEND:-0}" != "1" ]; then
24+
echo "NVTE_ENABLE_CUTEDSL_BACKEND is not 1; L0_cppunittest should have validated the CUDA path already."
25+
exit 0
26+
fi
27+
28+
# Override this env var to force the test to take the CUDA implementation path
29+
export NVTE_ENABLE_CUTEDSL_BACKEND=0
30+
31+
cd $TE_PATH/tests/cpp
32+
cmake -GNinja -Bbuild .
33+
cmake --build build
34+
export OMP_NUM_THREADS=$((NUM_PHYSICAL_CORES / NUM_PARALLEL_JOBS))
35+
ctest --test-dir build -j$NUM_PARALLEL_JOBS --output-junit $XML_LOG_DIR/ctest_cppunittest.xml

‎setup.py‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -64,6 +64,15 @@ def setup_common_extension() -> CMakeExtension:
6464
"-DCMAKE_CUDA_ARCHITECTURES={}".format(archs),
6565
f"-DCUDNN_FRONTEND_INCLUDE_DIR={cudnn_frontend_include_path()}",
6666
]
67+
68+
with_cutedsl = bool(int(os.getenv("NVTE_WITH_CUTEDSL", "0")))
69+
if with_cutedsl:
70+
tvm_ffi_include_dir = metadata.distribution("apache-tvm-ffi").locate_file("tvm_ffi/include")
71+
cmake_flags.append(f"-DTVM_FFI_INCLUDE_DIR={tvm_ffi_include_dir}")
72+
cmake_flags.append("-DNVTE_WITH_CUTEDSL=ON")
73+
else:
74+
cmake_flags.append("-DNVTE_WITH_CUTEDSL=OFF")
75+
6776
if bool(int(os.getenv("NVTE_UB_WITH_MPI", "0"))):
6877
assert (
6978
os.getenv("MPI_HOME") is not None

‎tests/cpp/CMakeLists.txt‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,5 +43,23 @@ include_directories(${CMAKE_SOURCE_DIR})
4343
find_package(CUDAToolkit REQUIRED)
4444
include(${CMAKE_SOURCE_DIR}/../../transformer_engine/common/cmake/cuDNN.cmake)
4545

46+
# Must match how libtransformer_engine.so was built, hence the env var default.
47+
option(NVTE_WITH_CUTEDSL
48+
"Link the tests against libpython, required when libtransformer_engine.so was built with NVTE_WITH_CUTEDSL=ON"
49+
"$ENV{NVTE_WITH_CUTEDSL}")
50+
4651
add_subdirectory(operator)
4752
add_subdirectory(util)
53+
54+
if(NVTE_WITH_CUTEDSL)
55+
# Executables that link libtransformerengine.so must also link to the Python::Python to provide
56+
# the definitions of these CPython symbols at runtime via dynamic linking.
57+
# Find the interpreter too so older CMake versions use its version to locate libpython.
58+
find_package(Python COMPONENTS Interpreter Development.Embed REQUIRED)
59+
foreach(test_target test_operator test_util)
60+
target_link_libraries(${test_target} PRIVATE
61+
"-Wl,--no-as-needed"
62+
Python::Python
63+
"-Wl,--as-needed")
64+
endforeach()
65+
endif()

‎tests/cpp/operator/CMakeLists.txt‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -52,5 +52,7 @@ list(APPEND test_operator_LINKER_LIBS CUDA::cudart GTest::gtest_main ${TE_LIB} C
5252
target_link_libraries(test_operator PUBLIC ${test_operator_LINKER_LIBS} OpenMP::OpenMP_CXX)
5353
target_compile_options(test_operator PRIVATE -O2 -fopenmp)
5454

55-
include(GoogleTest)
56-
gtest_discover_tests(test_operator DISCOVERY_TIMEOUT 600)
55+
# Keep all operator tests in one process so CUDA initialization and, when enabled,
56+
# the embedded Python interpreter and CuTeDSL kernel cache are initialized only once.
57+
# TODO(kainingz): This will only use 1 CPU and 1 GPU. Make it able to utilize all CPUs and GPUs later
58+
add_test(NAME test_operator COMMAND test_operator)

0 commit comments

Comments
 (0)