structuredAttrs is enabled Sourcing python-remove-tests-dir-hook Sourcing python-catch-conflicts-hook.sh Sourcing python-remove-bin-bytecode-hook.sh Sourcing pypa-build-hook Using pypaBuildPhase Sourcing python-runtime-deps-check-hook Using pythonRuntimeDepsCheckHook Sourcing pypa-install-hook Using pypaInstallPhase Sourcing python-imports-check-hook.sh Using pythonImportsCheckPhase Sourcing python-metadata-check-hook.sh Using pythonMetadataCheckPhase Sourcing python-namespaces-hook Sourcing python-catch-conflicts-hook.sh Sourcing setup-cuda-hook Sourcing pytest-check-hook Using pytestCheckPhase Using checkPhaseThreadLimitHook source: sourcing removeStubsFromRunpathHook.bash (hostOffset=0) (targetOffset=1) source: added removeStubsFromRunpathHookRegistration to prePhases Sourcing fix-elf-files.sh Running phase: removeStubsFromRunpathHookRegistration removeStubsFromRunpathHookRegistration: added removeStubsFromRunpath to postFixupHooks Running phase: unpackPhase unpacking source archive /nix/store/wff5yymdwhnqc7hx5lcxp9hf5yvd6i89-source source root is source setting SOURCE_DATE_EPOCH to timestamp 315619200 of file "source/xformers/utils.py" Running phase: patchPhase applying patch /nix/store/ai1qwfaq5vilgyl1sj9vxs1434ax99fp-xformers-fix-tests-pytorch-2.13.patch patching file CHANGELOG.md patching file xformers/checkpoint.py Running phase: updateAutotoolsGnuConfigScriptsPhase Running phase: configurePhase Executing setupCUDAToolkitCompilers no configure script, doing nothing Running phase: buildPhase Executing pypaBuildPhase Creating a wheel... pypa build flags: --no-isolation --outdir dist/ --wheel * Getting build dependencies for wheel... /build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry Error in cpuinfo: failed to parse the list of possible processors in /sys/devices/system/cpu/possible Error in cpuinfo: failed to parse the list of present processors in /sys/devices/system/cpu/present Error in cpuinfo: failed to parse both lists of possible and present processors W0930 10:42:57.984000 114 torch/utils/cpp_extension.py:178] No CUDA runtime is found, using CUDA_HOME='/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86' /nix/store/wf83qi4vq4ag4xl0gfni7m5sv2pzgfd5-python3.13-wheel-0.47.0/lib/python3.13/site-packages/wheel/bdist_wheel.py:4: FutureWarning: The 'wheel' package is no longer the canonical location of the 'bdist_wheel' command, and will be removed in a future release. Please update to setuptools v70.1 or later which contains an integrated version of this command. warn( /nix/store/8p92j2ppflf7igibazxiwsh3wa8y528w-python3.13-setuptools-83.0.0/lib/python3.13/site-packages/setuptools/dist.py:765: SetuptoolsDeprecationWarning: License classifiers are deprecated. !! ******************************************************************************** Please consider removing the following classifiers in favor of a SPDX license expression: License :: OSI Approved :: BSD License See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details. ******************************************************************************** !! self._finalize_license_expression() running egg_info creating xformers.egg-info writing xformers.egg-info/PKG-INFO writing dependency_links to xformers.egg-info/dependency_links.txt writing requirements to xformers.egg-info/requires.txt writing top-level names to xformers.egg-info/top_level.txt writing manifest file 'xformers.egg-info/SOURCES.txt' W0930 10:42:58.139000 114 torch/utils/cpp_extension.py:781] Attempted to use ninja as the BuildExtension backend but we could not find ninja.. Falling back to using the slow distutils backend. reading manifest file 'xformers.egg-info/SOURCES.txt' reading manifest template 'MANIFEST.in' adding license file 'LICENSE' writing manifest file 'xformers.egg-info/SOURCES.txt' * Building wheel... /build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry Error in cpuinfo: failed to parse the list of possible processors in /sys/devices/system/cpu/possible Error in cpuinfo: failed to parse the list of present processors in /sys/devices/system/cpu/present Error in cpuinfo: failed to parse both lists of possible and present processors W0930 10:43:00.298000 117 torch/utils/cpp_extension.py:178] No CUDA runtime is found, using CUDA_HOME='/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86' /nix/store/wf83qi4vq4ag4xl0gfni7m5sv2pzgfd5-python3.13-wheel-0.47.0/lib/python3.13/site-packages/wheel/bdist_wheel.py:4: FutureWarning: The 'wheel' package is no longer the canonical location of the 'bdist_wheel' command, and will be removed in a future release. Please update to setuptools v70.1 or later which contains an integrated version of this command. warn( /nix/store/8p92j2ppflf7igibazxiwsh3wa8y528w-python3.13-setuptools-83.0.0/lib/python3.13/site-packages/setuptools/dist.py:765: SetuptoolsDeprecationWarning: License classifiers are deprecated. !! ******************************************************************************** Please consider removing the following classifiers in favor of a SPDX license expression: License :: OSI Approved :: BSD License See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details. ******************************************************************************** !! self._finalize_license_expression() running bdist_wheel W0930 10:43:00.445000 117 torch/utils/cpp_extension.py:781] Attempted to use ninja as the BuildExtension backend but we could not find ninja.. Falling back to using the slow distutils backend. running build running build_py creating build/lib.linux-aarch64-cpython-313/xformers copying xformers/test.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/checkpoint.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/utils.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/attn_bias_utils.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/info.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/_deprecation_warning.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/_cpp_lib.py -> build/lib.linux-aarch64-cpython-313/xformers copying xformers/fwbw_overlap.py -> build/lib.linux-aarch64-cpython-313/xformers creating build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/sp24.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/rope_padded.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/rmsnorm.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/tiled_matmul.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/tree_attention.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/common.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/swiglu_op.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/seqpar.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/unbind.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/sequence_parallel_fused_ops.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/modpar_layers.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/indexing.py -> build/lib.linux-aarch64-cpython-313/xformers/ops copying xformers/ops/differentiable_collectives.py -> build/lib.linux-aarch64-cpython-313/xformers/ops creating build/lib.linux-aarch64-cpython-313/xformers/flash_attn_3 copying xformers/flash_attn_3/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/flash_attn_3 creating build/lib.linux-aarch64-cpython-313/xformers/components copying xformers/components/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/components creating build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_mem_eff_attention.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_tiled_matmul.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_sp24.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_indexing.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_merge_attentions.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_sequence_parallel_fused.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/benchmark_attn_decoding.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks copying xformers/benchmarks/utils.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks creating build/lib.linux-aarch64-cpython-313/xformers/sparse copying xformers/sparse/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/sparse copying xformers/sparse/blocksparse_tensor.py -> build/lib.linux-aarch64-cpython-313/xformers/sparse copying xformers/sparse/utils.py -> build/lib.linux-aarch64-cpython-313/xformers/sparse creating build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/profiler_dcgm_impl.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/profiler.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/api.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/profile_analyzer.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/profiler_dcgm.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/device_limits.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler copying xformers/profiler/find_slowest.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler creating build/lib.linux-aarch64-cpython-313/xformers/triton copying xformers/triton/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/triton copying xformers/triton/vararg_kernel.py -> build/lib.linux-aarch64-cpython-313/xformers/triton copying xformers/triton/importing.py -> build/lib.linux-aarch64-cpython-313/xformers/triton creating build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/torch_attention_compat.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/cutlass_blackwell.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/attn_bias.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/dispatch.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/merge_training.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/flash.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/flash3.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/ck.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/triton_splitk.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/cutlass.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/ck_splitk.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha copying xformers/ops/fmha/common.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha creating build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/k_index_select_cat.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/rope_padded_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/rmsnorm_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/matmul_perf_model.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/tiled_matmul_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton copying xformers/ops/_triton/k_scaled_index_add.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton creating build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton copying xformers/ops/fmha/_triton/splitk_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton copying xformers/ops/fmha/_triton/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton creating build/lib.linux-aarch64-cpython-313/xformers/components/attention copying xformers/components/attention/attention_patterns.py -> build/lib.linux-aarch64-cpython-313/xformers/components/attention copying xformers/components/attention/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/components/attention running build_ext W0930 10:43:00.504000 117 torch/utils/cpp_extension.py:657] There are no g++ version bounds defined for CUDA version 13.2 building 'xformers._C' extension creating build/temp.linux-aarch64-cpython-313/xformers/csrc/attention creating build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24 g++ -fno-strict-overflow -Wsign-compare -DNDEBUG -g -O3 -Wall -fPIC -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/attention/attention.cpp -o build/temp.linux-aarch64-cpython-313/xformers/csrc/attention/attention.o -O3 -std=c++17 -fopenmp -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/nvcc_info.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/nvcc_info.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 2.433 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 1.861 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/pt_stable_utils.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/pt_stable_utils.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 2.134 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 1.854 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/gemm.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/gemm.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 48 bytes gmem ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_20TileIteratorTensorOpIS1H_S1K_fS13_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_20TileIteratorTensorOpIS1H_S1K_fS13_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 214 registers, used 1 barriers ptxas info : Compile time = 171.988 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_20TileIteratorTensorOpIS1H_S1K_fSX_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_20TileIteratorTensorOpIS1H_S1K_fSX_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 220 registers, used 1 barriers ptxas info : Compile time = 172.818 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 218 registers, used 1 barriers ptxas info : Compile time = 164.534 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 224 registers, used 1 barriers ptxas info : Compile time = 163.473 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 214 registers, used 1 barriers ptxas info : Compile time = 170.488 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 220 registers, used 1 barriers ptxas info : Compile time = 174.017 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 218 registers, used 1 barriers ptxas info : Compile time = 169.267 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 226 registers, used 1 barriers ptxas info : Compile time = 166.935 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.700 ms ptxas info : 48 bytes gmem ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_20TileIteratorTensorOpIS1H_S1K_fS13_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_20TileIteratorTensorOpIS1H_S1K_fS13_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 214 registers, used 1 barriers ptxas info : Compile time = 172.816 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_20TileIteratorTensorOpIS1H_S1K_fSX_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_20TileIteratorTensorOpIS1H_S1K_fSX_EENS1W_18SharedLoadIteratorINS23_18CompactedThreadMapEfLi32EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 220 registers, used 1 barriers ptxas info : Compile time = 173.837 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 218 registers, used 1 barriers ptxas info : Compile time = 165.266 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_10bfloat16_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_20TileIteratorTensorOpIS1I_S1L_fSF_EENS1X_18SharedLoadIteratorINS24_18CompactedThreadMapEfLi32EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi1ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 224 registers, used 1 barriers ptxas info : Compile time = 163.301 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi64ELi128EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ESY_Li16EEELSV_1EfNSE_8RowMajorENSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S11_fS13_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_S13_SD_SF_fS13_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEES13_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 214 registers, used 1 barriers ptxas info : Compile time = 171.595 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout11ColumnMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi256ELi32EEELi256ENSH_ILi8ELi4EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_40ColumnMajorTensorOpMultiplicandCongruousILi16ELi64EEELi1ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_8RowMajorELi0ENSG_INSH_ILi128ELi64EEELi256ESJ_Li8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSX_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSF_Li0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSX_NS1F_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SX_SD_SF_fSX_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1S_S1S_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1R_Li1ENS1W_22PredicatedTileIteratorINS1W_26OutputTileOptimalThreadMapINS1W_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS20_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1V_4warp24FragmentIteratorTensorOpIS1H_S1K_fNSL_IfLi4ELb1EEESX_EENS25_25TileIteratorTensorOpMixedIS1H_S1K_fLi32ELi16ELi8ELi8ELb0EEENS1W_23SharedLoadIteratorMixedINS23_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1V_6thread17LinearCombinationISD_Li8EffLNS2E_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 220 registers, used 1 barriers ptxas info : Compile time = 175.636 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_NSE_11ColumnMajorELi0ENSG_INSH_ILi64ELi128EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_40ColumnMajorTensorOpMultiplicandCrosswiseILi16ELi64EEELi1ES10_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtSX_Li0ES1A_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S13_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_SX_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 218 registers, used 1 barriers ptxas info : Compile time = 169.881 ms ptxas info : Compiling entry function '_ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass6KernelINS_4gemm6kernel10SparseGemmINS1_11threadblock19SparseMmaMultistageINS1_9GemmShapeILi256ELi128ELi64EEENS_9transform11threadblock28PredicatedTileAccessIteratorINS_11MatrixShapeILi256ELi32EEENS_6half_tENS_6layout8RowMajorELi1ENS8_29PitchLinearWarpRakedThreadMapINS_16PitchLinearShapeILi32ELi256EEELi256ENSH_ILi4ELi8EEELi8EEENS_5ArrayISD_Li8ELb0EEELb0ENSE_9NoPermuteEEENS9_25RegularTileAccessIteratorISC_SD_NSE_37RowMajorTensorOpMultiplicandCrosswiseILi16ELi32EEELi0ESK_Li16EEELNS_4arch14CacheOperation4KindE1ENSA_INSB_ILi64ELi128EEESD_SF_Li0ENSG_INSH_ILi128ELi64EEELi256ENSH_ILi8ELi4EEELi8EEESM_Lb0ESN_EENSP_ISW_SD_NSE_37RowMajorTensorOpMultiplicandCongruousILi16ELi64EEELi0ESZ_Li16EEELSV_1EfSF_NSA_INSB_ILi256ELi4EEEtNSE_22ColumnMajorInterleavedILi2EEELi1ENS8_30PitchLinearStripminedThreadMapINSH_ILi512ELi2EEELi128ELi8EEENSL_ItLi8ELb0EEELb0ESN_EENSP_INSB_ILi512ELi2EEEtNSE_11ColumnMajorELi0ES19_Li16EEELSV_1ENS4_15SparseMmaPolicyINS1_4warp17SparseMmaTensorOpINS6_ILi64ELi64ELi64EEESD_SR_SD_S12_fSF_NS1G_17MmaTensorOpPolicyINST_9SparseMmaINS6_ILi16ELi8ELi32EEELi32ESD_SF_SD_S1D_fSF_NST_13OpMultiplyAddELNST_12SPFormatType4KindE0EEENSB_ILi1ELi1EEEEELi1ELb0EbEENSB_ILi0ELi0EEES1T_S1T_Li1EEELi4EbEENS_8epilogue11threadblock8EpilogueIS7_S1S_Li1ENS1X_22PredicatedTileIteratorINS1X_26OutputTileOptimalThreadMapINS1X_15OutputTileShapeILi128ELi8ELi4ELi1ELi1EEENS21_ILi1ELi8ELi1ELi1ELi8EEELi256ELi8ELi16EEESD_Lb0ESN_Lb0EEENS1W_4warp24FragmentIteratorTensorOpIS1I_S1L_fNSL_IfLi4ELb1EEESF_EENS26_25TileIteratorTensorOpMixedIS1I_S1L_fLi32ELi16ELi8ELi8ELb0EEENS1X_23SharedLoadIteratorMixedINS24_18CompactedThreadMapEfLi32ELi16ELi8ELi8ELb0EEENS1W_6thread17LinearCombinationISD_Li8EffLNS2F_9ScaleType4KindE0ELNS_15FloatRoundStyleE2ESD_EENSB_ILi0ELi8EEELi2ELi1EEENS4_30GemmIdentityThreadblockSwizzleILi3EEELb0EEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 226 registers, used 1 barriers ptxas info : Compile time = 168.442 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.771 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/meta_utils.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/meta_utils.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 48 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 2.116 ms ptxas info : 48 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 2.793 ms g++ -fno-strict-overflow -Wsign-compare -DNDEBUG -g -O3 -Wall -fPIC -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24.cpp -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24.o -O3 -std=c++17 -fopenmp -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24_apply.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_apply.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 52.814 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 50.078 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.588 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 54.318 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 49.700 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.577 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24_apply_dense_output.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_apply_dense_output.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb0ELb1EEEvNS_6ParamsIT_EE' for 'sm_120' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb0ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 80 registers, used 0 barriers ptxas info : Compile time = 39.943 ms ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb1ELb1EEEvNS_6ParamsIT_EE' for 'sm_120' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb1ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 59 registers, used 0 barriers ptxas info : Compile time = 45.983 ms ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb0ELb1EEEvNS_6ParamsIT_EE' for 'sm_120' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb0ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 80 registers, used 0 barriers ptxas info : Compile time = 38.088 ms ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb1ELb1EEEvNS_6ParamsIT_EE' for 'sm_120' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb1ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 59 registers, used 0 barriers ptxas info : Compile time = 46.816 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.636 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb0ELb1EEEvNS_6ParamsIT_EE' for 'sm_121' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb0ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 80 registers, used 0 barriers ptxas info : Compile time = 39.949 ms ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb1ELb1EEEvNS_6ParamsIT_EE' for 'sm_121' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb1ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 59 registers, used 0 barriers ptxas info : Compile time = 45.869 ms ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb0ELb1EEEvNS_6ParamsIT_EE' for 'sm_121' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb0ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 80 registers, used 0 barriers ptxas info : Compile time = 37.928 ms ptxas info : Compiling entry function '_ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb1ELb1EEEvNS_6ParamsIT_EE' for 'sm_121' ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb1ELb1EEEvNS_6ParamsIT_EE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 59 registers, used 0 barriers ptxas info : Compile time = 46.791 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.620 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24_gemm_sm90.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_gemm_sm90.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc xformers/csrc/sparse24/sparse24_gemm_sm90.cu(139): warning #177-D: member "::SparseRowwiseKernel::kElementOutAt" was declared but never referenced static constexpr auto kElementOutAt = torch::headeronly::ScalarType::BFloat16; ^ Remark: The warnings can be suppressed with "-diag-suppress " xformers/csrc/sparse24/sparse24_gemm_sm90.cu(139): warning #177-D: member "::SparseRowwiseKernel::kElementOutAt" was declared but never referenced static constexpr auto kElementOutAt = torch::headeronly::ScalarType::BFloat16; ^ Remark: The warnings can be suppressed with "-diag-suppress " ptxas info : 155 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 1.821 ms ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 1 barriers ptxas info : Compile time = 109.399 ms ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 162 registers, used 1 barriers ptxas info : Compile time = 240.021 ms ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalIN4cute5tupleIJiiiiEEENS1_10collective13CollectiveMmaINS1_40MainloopSm90TmaGmmaWarpSpecializedSparseILi4ENS5_IJNS4_1CILi2EEENSA_ILi1EEESC_EEENS1_47KernelTmaWarpSpecializedCooperativeFP8FastAccumEEENS5_IJNSA_ILi128EEENSA_ILi256EEESG_EEENS_12float_e4m3_tENS5_IJNS4_6LayoutINS5_IJiNS5_IJSB_iEEEiEEENS5_IJlNS5_IJSC_SB_EEElEEEEENSK_INS5_IJNS5_IJNSA_ILi64EEEiEEENS5_IJSG_iEEEiEEENS5_IJNS5_IJSG_NSA_ILi8192EEEEEENS5_IJSC_lEEElEEEEEEEESJ_NS5_IJlSC_lEEENS4_8TiledMMAINS4_8MMA_AtomIJNS4_4SM904GMMA6SPARSE32GMMA_64x256x64_F32E4M3E4M3_SS_TNILNS14_7ScaleInE1ELS17_1ELNS14_9SparseSelE0EEEEEENSK_ISD_NS5_IJSC_NSA_ILi0EEES1B_EEEEENS5_IJNS4_10UnderscoreES1E_S1E_EEEEENS4_13SM90_TMA_LOADENS4_14ComposedLayoutINS4_7SwizzleILi2ELi4ELi3EEENS4_25smem_sparse_ptr_flag_bitsILi2ELi8EEENSK_INS5_IJNS5_IJSC_NSA_ILi8EEEEEENS5_IJSB_SQ_EEEEEENS5_IJNS5_IJS1B_SG_EEESN_EEEEEEEvNS4_8identityENS4_23SM90_TMA_LOAD_MULTICASTENS1I_INS1J_ILi3ELi4ELi3EEENS4_18smem_ptr_flag_bitsILi8EEENSK_INS5_IJS1N_SG_EEENS5_IJSG_SC_EEEEEEEvS1V_EENS_8epilogue10collective18CollectiveEpilogueINS25_22Sm90TmaWarpSpecializedILi4ELi2ELi16ELb1ELb0EEEJSI_NS5_IJSG_NSA_ILi32EEEEEENS_10bfloat16_tES10_S2C_S10_NS25_6fusion15Sm90TreeVisitorINS2D_11Sm90ComputeIN58_GLOBAL__N__81c05b03_21_sparse24_gemm_sm90_cu_4cca563f_2808identityES2C_fLNS_15FloatRoundStyleE2EvEEJNS2E_INS2F_INS_10multipliesEffLS2I_2EvEEJNS2D_16Sm90RowBroadcastILi0ESI_ffNS5_IJS1B_SC_S1B_EEELi4ELb1EEENS2E_IS2L_JNS2D_16Sm90ColBroadcastILi0ESI_ffS1C_Li4ELb1EEENS2D_12Sm90AccFetchEEEEEEEEEES1H_NS1I_IS1K_NS1Y_ILi16EEENSK_INS5_IJS1N_S2A_EEENS5_IJS2A_SC_EEEEEEENS4_17SM75_U32x4_LDSM_NENS4_14SM90_TMA_STOREES2Z_NS4_17SM90_U32x4_STSM_NENS4_9Copy_AtomIJS32_NS_6half_tEEEEvEEENS1_19PersistentSchedulerEvEEEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalIN4cute5tupleIJiiiiEEENS1_10collective13CollectiveMmaINS1_40MainloopSm90TmaGmmaWarpSpecializedSparseILi4ENS5_IJNS4_1CILi2EEENSA_ILi1EEESC_EEENS1_47KernelTmaWarpSpecializedCooperativeFP8FastAccumEEENS5_IJNSA_ILi128EEENSA_ILi256EEESG_EEENS_12float_e4m3_tENS5_IJNS4_6LayoutINS5_IJiNS5_IJSB_iEEEiEEENS5_IJlNS5_IJSC_SB_EEElEEEEENSK_INS5_IJNS5_IJNSA_ILi64EEEiEEENS5_IJSG_iEEEiEEENS5_IJNS5_IJSG_NSA_ILi8192EEEEEENS5_IJSC_lEEElEEEEEEEESJ_NS5_IJlSC_lEEENS4_8TiledMMAINS4_8MMA_AtomIJNS4_4SM904GMMA6SPARSE32GMMA_64x256x64_F32E4M3E4M3_SS_TNILNS14_7ScaleInE1ELS17_1ELNS14_9SparseSelE0EEEEEENSK_ISD_NS5_IJSC_NSA_ILi0EEES1B_EEEEENS5_IJNS4_10UnderscoreES1E_S1E_EEEEENS4_13SM90_TMA_LOADENS4_14ComposedLayoutINS4_7SwizzleILi2ELi4ELi3EEENS4_25smem_sparse_ptr_flag_bitsILi2ELi8EEENSK_INS5_IJNS5_IJSC_NSA_ILi8EEEEEENS5_IJSB_SQ_EEEEEENS5_IJNS5_IJS1B_SG_EEESN_EEEEEEEvNS4_8identityENS4_23SM90_TMA_LOAD_MULTICASTENS1I_INS1J_ILi3ELi4ELi3EEENS4_18smem_ptr_flag_bitsILi8EEENSK_INS5_IJS1N_SG_EEENS5_IJSG_SC_EEEEEEEvS1V_EENS_8epilogue10collective18CollectiveEpilogueINS25_22Sm90TmaWarpSpecializedILi4ELi2ELi16ELb1ELb0EEEJSI_NS5_IJSG_NSA_ILi32EEEEEENS_10bfloat16_tES10_S2C_S10_NS25_6fusion15Sm90TreeVisitorINS2D_11Sm90ComputeIN58_GLOBAL__N__81c05b03_21_sparse24_gemm_sm90_cu_4cca563f_2808identityES2C_fLNS_15FloatRoundStyleE2EvEEJNS2E_INS2F_INS_10multipliesEffLS2I_2EvEEJNS2D_16Sm90RowBroadcastILi0ESI_ffNS5_IJS1B_SC_S1B_EEELi4ELb1EEENS2E_IS2L_JNS2D_16Sm90ColBroadcastILi0ESI_ffS1C_Li4ELb1EEENS2D_12Sm90AccFetchEEEEEEEEEES1H_NS1I_IS1K_NS1Y_ILi16EEENSK_INS5_IJS1N_S2A_EEENS5_IJS2A_SC_EEEEEEENS4_17SM75_U32x4_LDSM_NENS4_14SM90_TMA_STOREES2Z_NS4_17SM90_U32x4_STSM_NENS4_9Copy_AtomIJS32_NS_6half_tEEEEvEEENS1_19PersistentSchedulerEvEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 registers, used 0 barriers ptxas info : Compile time = 1.500 ms ptxas info : 155 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 1.639 ms ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 1 barriers ptxas info : Compile time = 109.581 ms ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 162 registers, used 1 barriers ptxas info : Compile time = 240.732 ms ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalIN4cute5tupleIJiiiiEEENS1_10collective13CollectiveMmaINS1_40MainloopSm90TmaGmmaWarpSpecializedSparseILi4ENS5_IJNS4_1CILi2EEENSA_ILi1EEESC_EEENS1_47KernelTmaWarpSpecializedCooperativeFP8FastAccumEEENS5_IJNSA_ILi128EEENSA_ILi256EEESG_EEENS_12float_e4m3_tENS5_IJNS4_6LayoutINS5_IJiNS5_IJSB_iEEEiEEENS5_IJlNS5_IJSC_SB_EEElEEEEENSK_INS5_IJNS5_IJNSA_ILi64EEEiEEENS5_IJSG_iEEEiEEENS5_IJNS5_IJSG_NSA_ILi8192EEEEEENS5_IJSC_lEEElEEEEEEEESJ_NS5_IJlSC_lEEENS4_8TiledMMAINS4_8MMA_AtomIJNS4_4SM904GMMA6SPARSE32GMMA_64x256x64_F32E4M3E4M3_SS_TNILNS14_7ScaleInE1ELS17_1ELNS14_9SparseSelE0EEEEEENSK_ISD_NS5_IJSC_NSA_ILi0EEES1B_EEEEENS5_IJNS4_10UnderscoreES1E_S1E_EEEEENS4_13SM90_TMA_LOADENS4_14ComposedLayoutINS4_7SwizzleILi2ELi4ELi3EEENS4_25smem_sparse_ptr_flag_bitsILi2ELi8EEENSK_INS5_IJNS5_IJSC_NSA_ILi8EEEEEENS5_IJSB_SQ_EEEEEENS5_IJNS5_IJS1B_SG_EEESN_EEEEEEEvNS4_8identityENS4_23SM90_TMA_LOAD_MULTICASTENS1I_INS1J_ILi3ELi4ELi3EEENS4_18smem_ptr_flag_bitsILi8EEENSK_INS5_IJS1N_SG_EEENS5_IJSG_SC_EEEEEEEvS1V_EENS_8epilogue10collective18CollectiveEpilogueINS25_22Sm90TmaWarpSpecializedILi4ELi2ELi16ELb1ELb0EEEJSI_NS5_IJSG_NSA_ILi32EEEEEENS_10bfloat16_tES10_S2C_S10_NS25_6fusion15Sm90TreeVisitorINS2D_11Sm90ComputeIN58_GLOBAL__N__81c05b03_21_sparse24_gemm_sm90_cu_4cca563f_2808identityES2C_fLNS_15FloatRoundStyleE2EvEEJNS2E_INS2F_INS_10multipliesEffLS2I_2EvEEJNS2D_16Sm90RowBroadcastILi0ESI_ffNS5_IJS1B_SC_S1B_EEELi4ELb1EEENS2E_IS2L_JNS2D_16Sm90ColBroadcastILi0ESI_ffS1C_Li4ELb1EEENS2D_12Sm90AccFetchEEEEEEEEEES1H_NS1I_IS1K_NS1Y_ILi16EEENSK_INS5_IJS1N_S2A_EEENS5_IJS2A_SC_EEEEEEENS4_17SM75_U32x4_LDSM_NENS4_14SM90_TMA_STOREES2Z_NS4_17SM90_U32x4_STSM_NENS4_9Copy_AtomIJS32_NS_6half_tEEEEvEEENS1_19PersistentSchedulerEvEEEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalIN4cute5tupleIJiiiiEEENS1_10collective13CollectiveMmaINS1_40MainloopSm90TmaGmmaWarpSpecializedSparseILi4ENS5_IJNS4_1CILi2EEENSA_ILi1EEESC_EEENS1_47KernelTmaWarpSpecializedCooperativeFP8FastAccumEEENS5_IJNSA_ILi128EEENSA_ILi256EEESG_EEENS_12float_e4m3_tENS5_IJNS4_6LayoutINS5_IJiNS5_IJSB_iEEEiEEENS5_IJlNS5_IJSC_SB_EEElEEEEENSK_INS5_IJNS5_IJNSA_ILi64EEEiEEENS5_IJSG_iEEEiEEENS5_IJNS5_IJSG_NSA_ILi8192EEEEEENS5_IJSC_lEEElEEEEEEEESJ_NS5_IJlSC_lEEENS4_8TiledMMAINS4_8MMA_AtomIJNS4_4SM904GMMA6SPARSE32GMMA_64x256x64_F32E4M3E4M3_SS_TNILNS14_7ScaleInE1ELS17_1ELNS14_9SparseSelE0EEEEEENSK_ISD_NS5_IJSC_NSA_ILi0EEES1B_EEEEENS5_IJNS4_10UnderscoreES1E_S1E_EEEEENS4_13SM90_TMA_LOADENS4_14ComposedLayoutINS4_7SwizzleILi2ELi4ELi3EEENS4_25smem_sparse_ptr_flag_bitsILi2ELi8EEENSK_INS5_IJNS5_IJSC_NSA_ILi8EEEEEENS5_IJSB_SQ_EEEEEENS5_IJNS5_IJS1B_SG_EEESN_EEEEEEEvNS4_8identityENS4_23SM90_TMA_LOAD_MULTICASTENS1I_INS1J_ILi3ELi4ELi3EEENS4_18smem_ptr_flag_bitsILi8EEENSK_INS5_IJS1N_SG_EEENS5_IJSG_SC_EEEEEEEvS1V_EENS_8epilogue10collective18CollectiveEpilogueINS25_22Sm90TmaWarpSpecializedILi4ELi2ELi16ELb1ELb0EEEJSI_NS5_IJSG_NSA_ILi32EEEEEENS_10bfloat16_tES10_S2C_S10_NS25_6fusion15Sm90TreeVisitorINS2D_11Sm90ComputeIN58_GLOBAL__N__81c05b03_21_sparse24_gemm_sm90_cu_4cca563f_2808identityES2C_fLNS_15FloatRoundStyleE2EvEEJNS2E_INS2F_INS_10multipliesEffLS2I_2EvEEJNS2D_16Sm90RowBroadcastILi0ESI_ffNS5_IJS1B_SC_S1B_EEELi4ELb1EEENS2E_IS2L_JNS2D_16Sm90ColBroadcastILi0ESI_ffS1C_Li4ELb1EEENS2D_12Sm90AccFetchEEEEEEEEEES1H_NS1I_IS1K_NS1Y_ILi16EEENSK_INS5_IJS1N_S2A_EEENS5_IJS2A_SC_EEEEEEENS4_17SM75_U32x4_LDSM_NENS4_14SM90_TMA_STOREES2Z_NS4_17SM90_U32x4_STSM_NENS4_9Copy_AtomIJS32_NS_6half_tEEEEvEEENS1_19PersistentSchedulerEvEEEEvNT_6ParamsE 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 registers, used 0 barriers ptxas info : Compile time = 1.556 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24_largest_mask_2d.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_largest_mask_2d.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 1.753 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 134.253 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 136.809 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 51.122 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE' for 'sm_120' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 54.085 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 1.708 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 135.907 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 138.598 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 51.033 ms ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE' for 'sm_121' ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size ptxas info : Compile time = 54.259 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24_pack.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_pack.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 314.039 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.251 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.670 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 156.318 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 293.931 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 309.011 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.945 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.869 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 154.925 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 304.721 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 310.663 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 124.588 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 125.155 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 153.806 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 292.594 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 306.241 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.673 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.637 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 154.419 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_120' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 301.579 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.784 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 315.079 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.593 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 127.147 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 156.576 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 294.921 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 309.906 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.965 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.687 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 155.787 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 306.276 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 310.619 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 124.817 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 125.454 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 154.234 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 292.750 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 306.190 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 126.516 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_ 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size ptxas info : Compile time = 127.104 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 96 registers, used 0 barriers ptxas info : Compile time = 154.725 ms ptxas info : Compiling entry function '_ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_' for 'sm_121' ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_ 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size ptxas info : Compile time = 303.476 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.791 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparse24_pack_test.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_pack_test.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN58_GLOBAL__N__0b8b56aa_21_sparse24_pack_test_cu_4cca563f_34924meta_shuffle_test_kernelEN5torch10headeronly6detail27GenericPackedTensorAccessorINS2_14TensorAccessorIN3c1018HeaderOnlyArrayRefIlEElLm2ENS1_16DefaultPtrTraitsElEENS2_26HeaderOnlyIndexBoundsCheckILm3ElEElLm3ES8_lEESC_b' for 'sm_121' ptxas info : Function properties for _ZN58_GLOBAL__N__0b8b56aa_21_sparse24_pack_test_cu_4cca563f_34924meta_shuffle_test_kernelEN5torch10headeronly6detail27GenericPackedTensorAccessorINS2_14TensorAccessorIN3c1018HeaderOnlyArrayRefIlEElLm2ENS1_16DefaultPtrTraitsElEENS2_26HeaderOnlyIndexBoundsCheckILm3ElEElLm3ES8_lEESC_b 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 26 registers, used 0 barriers ptxas info : Compile time = 10.000 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.886 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN58_GLOBAL__N__0b8b56aa_21_sparse24_pack_test_cu_4cca563f_34924meta_shuffle_test_kernelEN5torch10headeronly6detail27GenericPackedTensorAccessorINS2_14TensorAccessorIN3c1018HeaderOnlyArrayRefIlEElLm2ENS1_16DefaultPtrTraitsElEENS2_26HeaderOnlyIndexBoundsCheckILm3ElEElLm3ES8_lEESC_b' for 'sm_120' ptxas info : Function properties for _ZN58_GLOBAL__N__0b8b56aa_21_sparse24_pack_test_cu_4cca563f_34924meta_shuffle_test_kernelEN5torch10headeronly6detail27GenericPackedTensorAccessorINS2_14TensorAccessorIN3c1018HeaderOnlyArrayRefIlEElLm2ENS1_16DefaultPtrTraitsElEENS2_26HeaderOnlyIndexBoundsCheckILm3ElEElLm3ES8_lEESC_b 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 26 registers, used 0 barriers ptxas info : Compile time = 9.690 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.894 ms /nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/bin/nvcc -I/build/source/xformers/csrc -I/build/source/third_party/cutlass/include -I/build/source/third_party/cutlass/tools/util/include -I/build/source/third_party/cutlass/examples -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include -I/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/include/torch/csrc/api/include -I/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/include -I/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/include/python3.13 -c xformers/csrc/sparse24/sparseNM_dense.cu -o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparseNM_dense.o -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr --compiler-options '-fPIC' -DHAS_PYTORCH --use_fast_math -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ --extended-lambda -D_ENABLE_EXTENDED_ALIGNED_STORAGE -std=c++17 --generate-line-info -DNDEBUG --threads 4 --ptxas-options=-v -DTORCH_STABLE_ONLY -DTORCH_TARGET_VERSION=0x020a000000000000 -DUSE_CUDA --ptxas-options=-O2 --ptxas-options=-allow-expensive-optimizations=true -DTORCH_API_INCLUDE_EXTENSION_H -DTORCH_EXTENSION_NAME=_C -gencode=arch=compute_120,code=sm_120 -gencode=arch=compute_121,code=sm_121 -ccbin gcc ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 registers, used 0 barriers ptxas info : Compile time = 84.550 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 23 registers, used 0 barriers ptxas info : Compile time = 71.719 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 32 registers, used 0 barriers ptxas info : Compile time = 87.689 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 28 registers, used 0 barriers ptxas info : Compile time = 77.425 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 registers, used 0 barriers ptxas info : Compile time = 81.265 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 23 registers, used 0 barriers ptxas info : Compile time = 68.931 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 32 registers, used 0 barriers ptxas info : Compile time = 85.886 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_120' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 28 registers, used 0 barriers ptxas info : Compile time = 73.455 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.676 ms ptxas info : 46 bytes gmem ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 registers, used 0 barriers ptxas info : Compile time = 84.736 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 23 registers, used 0 barriers ptxas info : Compile time = 71.683 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 32 registers, used 0 barriers ptxas info : Compile time = 87.624 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb1EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 28 registers, used 0 barriers ptxas info : Compile time = 77.222 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 24 registers, used 0 barriers ptxas info : Compile time = 81.803 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi4EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 23 registers, used 0 barriers ptxas info : Compile time = 68.885 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E0_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 32 registers, used 0 barriers ptxas info : Compile time = 85.740 ms ptxas info : Compiling entry function '_ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_' for 'sm_121' ptxas info : Function properties for _ZN54_GLOBAL__N__11a0e9b9_17_sparseNM_dense_cu_4cca563f_37221sparseNM_dense_kernelILi2ELi8EN7cutlass10bfloat16_tEZNS_14sparseNM_denseILb0EEEN5torch6stable6TensorES6_NSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEllEUlT_E_EEvPKT1_lPSF_lT2_ 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 28 registers, used 0 barriers ptxas info : Compile time = 73.030 ms ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121' ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads ptxas info : Used 4 registers, used 0 barriers ptxas info : Compile time = 0.688 ms g++ -fno-strict-overflow -Wsign-compare -DNDEBUG -g -O3 -Wall -shared -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-bzip2-1.0.8/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-libffi-3.8.0/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-util-linux-minimal-2.42.3-lib/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-ncurses-6.6/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-xz-5.8.3/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-zlib-1.3.2/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-openssl-3.6.4/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-sqlite-3.53.3/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-mpdecimal-4.0.1/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-expat-2.8.4/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-bluez-headers-5.87/lib -L/nix/store/i8an849ir6g3f4n2wrsm88s2qgrzjayq-tzdata-2026c/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-gdbm-1.26-lib/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-readline-8.3p3/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-bzip2-1.0.8/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-libffi-3.8.0/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-util-linux-minimal-2.42.3-lib/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-ncurses-6.6/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-xz-5.8.3/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-zlib-1.3.2/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-openssl-3.6.4/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-sqlite-3.53.3/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-mpdecimal-4.0.1/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-expat-2.8.4/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-bluez-headers-5.87/lib -L/nix/store/i8an849ir6g3f4n2wrsm88s2qgrzjayq-tzdata-2026c/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-gdbm-1.26-lib/lib -L/nix/store/eeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee-readline-8.3p3/lib build/temp.linux-aarch64-cpython-313/xformers/csrc/attention/attention.o build/temp.linux-aarch64-cpython-313/xformers/csrc/nvcc_info.o build/temp.linux-aarch64-cpython-313/xformers/csrc/pt_stable_utils.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/gemm.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/meta_utils.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_apply.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_apply_dense_output.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_gemm_sm90.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_largest_mask_2d.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_pack.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparse24_pack_test.o build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24/sparseNM_dense.o -L/nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/lib -L/nix/store/hrjlgxvpfzawmq9ji5iaxyz53lyivlnf-cuda13.2-cuda_nvcc-13.2.86/lib64 -L/nix/store/s81lcrdqrwmvnqdl8h7m711xafjdx17f-python3-3.13.15/lib -lc10 -ltorch -ltorch_cpu -ltorch_python -lcudart -lc10_cuda -ltorch_cuda -o build/lib.linux-aarch64-cpython-313/xformers/_C.so installing to build/bdist.linux-aarch64/wheel running install running install_lib creating build/bdist.linux-aarch64/wheel creating build/bdist.linux-aarch64/wheel/xformers copying build/lib.linux-aarch64-cpython-313/xformers/test.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/_deprecation_warning.py -> build/bdist.linux-aarch64/wheel/./xformers creating build/bdist.linux-aarch64/wheel/xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_sp24.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_merge_attentions.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_attn_decoding.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_mem_eff_attention.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_tiled_matmul.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/utils.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_indexing.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_sequence_parallel_fused.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks copying build/lib.linux-aarch64-cpython-313/xformers/utils.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/checkpoint.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/fwbw_overlap.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/_C.so -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/_cpp_lib.py -> build/bdist.linux-aarch64/wheel/./xformers creating build/bdist.linux-aarch64/wheel/xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/rmsnorm.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/swiglu_op.py -> build/bdist.linux-aarch64/wheel/./xformers/ops creating build/bdist.linux-aarch64/wheel/xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/cutlass.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/triton_splitk.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha creating build/bdist.linux-aarch64/wheel/xformers/ops/fmha/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton/splitk_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/ck_splitk.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/ck.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/dispatch.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/attn_bias.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/flash3.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/flash.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/cutlass_blackwell.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/merge_training.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/common.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/torch_attention_compat.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha copying build/lib.linux-aarch64-cpython-313/xformers/ops/common.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/modpar_layers.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/tree_attention.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/indexing.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/unbind.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/differentiable_collectives.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/tiled_matmul.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/sp24.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/seqpar.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/rope_padded.py -> build/bdist.linux-aarch64/wheel/./xformers/ops copying build/lib.linux-aarch64-cpython-313/xformers/ops/sequence_parallel_fused_ops.py -> build/bdist.linux-aarch64/wheel/./xformers/ops creating build/bdist.linux-aarch64/wheel/xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/k_scaled_index_add.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/matmul_perf_model.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/rope_padded_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/rmsnorm_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/tiled_matmul_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/k_index_select_cat.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton creating build/bdist.linux-aarch64/wheel/xformers/flash_attn_3 copying build/lib.linux-aarch64-cpython-313/xformers/flash_attn_3/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/flash_attn_3 copying build/lib.linux-aarch64-cpython-313/xformers/info.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/cpp_lib.json -> build/bdist.linux-aarch64/wheel/./xformers creating build/bdist.linux-aarch64/wheel/xformers/triton copying build/lib.linux-aarch64-cpython-313/xformers/triton/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/triton copying build/lib.linux-aarch64-cpython-313/xformers/triton/importing.py -> build/bdist.linux-aarch64/wheel/./xformers/triton copying build/lib.linux-aarch64-cpython-313/xformers/triton/vararg_kernel.py -> build/bdist.linux-aarch64/wheel/./xformers/triton copying build/lib.linux-aarch64-cpython-313/xformers/attn_bias_utils.py -> build/bdist.linux-aarch64/wheel/./xformers copying build/lib.linux-aarch64-cpython-313/xformers/version.py -> build/bdist.linux-aarch64/wheel/./xformers creating build/bdist.linux-aarch64/wheel/xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profiler_dcgm_impl.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/find_slowest.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profiler.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/api.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profile_analyzer.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profiler_dcgm.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler copying build/lib.linux-aarch64-cpython-313/xformers/profiler/device_limits.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler creating build/bdist.linux-aarch64/wheel/xformers/components creating build/bdist.linux-aarch64/wheel/xformers/components/attention copying build/lib.linux-aarch64-cpython-313/xformers/components/attention/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/components/attention copying build/lib.linux-aarch64-cpython-313/xformers/components/attention/attention_patterns.py -> build/bdist.linux-aarch64/wheel/./xformers/components/attention copying build/lib.linux-aarch64-cpython-313/xformers/components/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/components creating build/bdist.linux-aarch64/wheel/xformers/sparse copying build/lib.linux-aarch64-cpython-313/xformers/sparse/utils.py -> build/bdist.linux-aarch64/wheel/./xformers/sparse copying build/lib.linux-aarch64-cpython-313/xformers/sparse/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/sparse copying build/lib.linux-aarch64-cpython-313/xformers/sparse/blocksparse_tensor.py -> build/bdist.linux-aarch64/wheel/./xformers/sparse running install_egg_info running egg_info writing xformers.egg-info/PKG-INFO writing dependency_links to xformers.egg-info/dependency_links.txt writing requirements to xformers.egg-info/requires.txt writing top-level names to xformers.egg-info/top_level.txt reading manifest file 'xformers.egg-info/SOURCES.txt' reading manifest template 'MANIFEST.in' adding license file 'LICENSE' writing manifest file 'xformers.egg-info/SOURCES.txt' Copying xformers.egg-info to build/bdist.linux-aarch64/wheel/./xformers-0.0.35-py3.13.egg-info running install_scripts creating build/bdist.linux-aarch64/wheel/xformers-0.0.35.dist-info/WHEEL creating '/build/source/dist/.tmp-8e8qf1el/xformers-0.0.35-py39-none-linux_aarch64.whl' and adding 'build/bdist.linux-aarch64/wheel' to it adding 'xformers/_C.so' adding 'xformers/__init__.py' adding 'xformers/_cpp_lib.py' adding 'xformers/_deprecation_warning.py' adding 'xformers/attn_bias_utils.py' adding 'xformers/checkpoint.py' adding 'xformers/cpp_lib.json' adding 'xformers/fwbw_overlap.py' adding 'xformers/info.py' adding 'xformers/test.py' adding 'xformers/utils.py' adding 'xformers/version.py' adding 'xformers/benchmarks/__init__.py' adding 'xformers/benchmarks/benchmark_attn_decoding.py' adding 'xformers/benchmarks/benchmark_indexing.py' adding 'xformers/benchmarks/benchmark_mem_eff_attention.py' adding 'xformers/benchmarks/benchmark_merge_attentions.py' adding 'xformers/benchmarks/benchmark_sequence_parallel_fused.py' adding 'xformers/benchmarks/benchmark_sp24.py' adding 'xformers/benchmarks/benchmark_tiled_matmul.py' adding 'xformers/benchmarks/utils.py' adding 'xformers/components/__init__.py' adding 'xformers/components/attention/__init__.py' adding 'xformers/components/attention/attention_patterns.py' adding 'xformers/flash_attn_3/__init__.py' adding 'xformers/ops/__init__.py' adding 'xformers/ops/common.py' adding 'xformers/ops/differentiable_collectives.py' adding 'xformers/ops/indexing.py' adding 'xformers/ops/modpar_layers.py' adding 'xformers/ops/rmsnorm.py' adding 'xformers/ops/rope_padded.py' adding 'xformers/ops/seqpar.py' adding 'xformers/ops/sequence_parallel_fused_ops.py' adding 'xformers/ops/sp24.py' adding 'xformers/ops/swiglu_op.py' adding 'xformers/ops/tiled_matmul.py' adding 'xformers/ops/tree_attention.py' adding 'xformers/ops/unbind.py' adding 'xformers/ops/_triton/__init__.py' adding 'xformers/ops/_triton/k_index_select_cat.py' adding 'xformers/ops/_triton/k_scaled_index_add.py' adding 'xformers/ops/_triton/matmul_perf_model.py' adding 'xformers/ops/_triton/rmsnorm_kernels.py' adding 'xformers/ops/_triton/rope_padded_kernels.py' adding 'xformers/ops/_triton/tiled_matmul_kernels.py' adding 'xformers/ops/fmha/__init__.py' adding 'xformers/ops/fmha/attn_bias.py' adding 'xformers/ops/fmha/ck.py' adding 'xformers/ops/fmha/ck_splitk.py' adding 'xformers/ops/fmha/common.py' adding 'xformers/ops/fmha/cutlass.py' adding 'xformers/ops/fmha/cutlass_blackwell.py' adding 'xformers/ops/fmha/dispatch.py' adding 'xformers/ops/fmha/flash.py' adding 'xformers/ops/fmha/flash3.py' adding 'xformers/ops/fmha/merge_training.py' adding 'xformers/ops/fmha/torch_attention_compat.py' adding 'xformers/ops/fmha/triton_splitk.py' adding 'xformers/ops/fmha/_triton/__init__.py' adding 'xformers/ops/fmha/_triton/splitk_kernels.py' adding 'xformers/profiler/__init__.py' adding 'xformers/profiler/api.py' adding 'xformers/profiler/device_limits.py' adding 'xformers/profiler/find_slowest.py' adding 'xformers/profiler/profile_analyzer.py' adding 'xformers/profiler/profiler.py' adding 'xformers/profiler/profiler_dcgm.py' adding 'xformers/profiler/profiler_dcgm_impl.py' adding 'xformers/sparse/__init__.py' adding 'xformers/sparse/blocksparse_tensor.py' adding 'xformers/sparse/utils.py' adding 'xformers/triton/__init__.py' adding 'xformers/triton/importing.py' adding 'xformers/triton/vararg_kernel.py' adding 'xformers-0.0.35.dist-info/licenses/LICENSE' adding 'xformers-0.0.35.dist-info/METADATA' adding 'xformers-0.0.35.dist-info/WHEEL' adding 'xformers-0.0.35.dist-info/top_level.txt' adding 'xformers-0.0.35.dist-info/RECORD' removing build/bdist.linux-aarch64/wheel Successfully built xformers-0.0.35-py39-none-linux_aarch64.whl Finished creating a wheel... Finished executing pypaBuildPhase buildPhase completed in 3 minutes 13 seconds Running phase: pythonRuntimeDepsCheckHook Executing pythonRuntimeDepsCheck Checking runtime dependencies for xformers-0.0.35-py39-none-linux_aarch64.whl Finished executing pythonRuntimeDepsCheck Running phase: installPhase Executing pypaInstallPhase Successfully installed xformers-0.0.35-py39-none-linux_aarch64.whl Finished executing pypaInstallPhase Running phase: pythonOutputDistPhase Executing pythonOutputDistPhase Finished executing pythonOutputDistPhase Running phase: fixupPhase shrinking RPATHs of ELF executables and libraries in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist checking for references to /build/ in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist... patching script interpreter paths in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist shrinking RPATHs of ELF executables and libraries in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist shrinking RPATHs of ELF executables and libraries in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35 shrinking /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib/python3.13/site-packages/xformers/_C.so checking for references to /build/ in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35... patching script interpreter paths in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35 stripping (with command strip and flags -S -p) in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib shrinking RPATHs of ELF executables and libraries in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35 shrinking /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib/python3.13/site-packages/xformers/_C.so Executing pythonRemoveTestsDir Finished executing pythonRemoveTestsDir autoFixElfFiles: using removeStubsFromRunpath to fix /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib/python3.13/site-packages/xformers/_C.so Running phase: installCheckPhase setting known thread variables to default: 1 OMP_NUM_THREADS=1 OPENBLAS_NUM_THREADS=1 MKL_NUM_THREADS=1 BLIS_NUM_THREADS=1 VECLIB_MAXIMUM_THREADS=1 NUMBA_NUM_THREADS=1 NUMEXPR_NUM_THREADS=1 no Makefile or custom installCheckPhase, doing nothing Running phase: pythonCatchConflictsPhase Running phase: pythonRemoveBinBytecodePhase Running phase: pythonImportsCheckPhase Executing pythonImportsCheckPhase Check whether the following modules can be imported: xformers xformers.components xformers.components.attention /build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry Running phase: pythonMetadataCheckPhase Executing pythonMetadataCheckPhase Running phase: pytestCheckPhase Executing pytestCheckPhase /build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry xFormers 0.0.35 memory_efficient_attention.ckF: unavailable memory_efficient_attention.ckB: unavailable memory_efficient_attention.ck_splitKF: unavailable memory_efficient_attention.cutlassF-pt: available memory_efficient_attention.cutlassB-pt: available memory_efficient_attention.cutlassF-blackwell: unavailable memory_efficient_attention.cutlassB-blackwell: unavailable memory_efficient_attention.fa2F@2.5.7-pt: available memory_efficient_attention.fa2B@2.5.7-pt: available memory_efficient_attention.fa3F@0.0.0: unavailable memory_efficient_attention.fa3B@0.0.0: unavailable memory_efficient_attention.fa3F_splitKV@0.0.0: unavailable memory_efficient_attention.triton_splitKF: available indexing.scaled_index_addF: unavailable indexing.scaled_index_addB: unavailable indexing.index_select: unavailable sp24.sparse24_sparsify_both_ways: available sp24.sparse24_apply: available sp24.sparse24_apply_dense_output: available sp24._sparse24_gemm: available sp24._cslt_sparse_mm_search@0.8.1: available sp24._cslt_sparse_mm@0.8.1: available is_triton_available: False pytorch.version: 2.13.0 pytorch.cuda: not available dcgm_profiler: unavailable build.info: available build.cuda_version: 1302 build.hip_version: None build.python_version: 3.13.15 build.torch_version: 2.13.0 build.env.TORCH_CUDA_ARCH_LIST: 12.0;12.1 build.env.PYTORCH_ROCM_ARCH: None build.env.XFORMERS_BUILD_TYPE: None build.env.XFORMERS_ENABLE_DEBUG_ASSERTIONS: None build.env.NVCC_FLAGS: None build.env.XFORMERS_PACKAGE_FROM: None build.nvcc_version: 13.2.86 source.privacy: open source setting known thread variables to default: 1 OMP_NUM_THREADS=1 OPENBLAS_NUM_THREADS=1 MKL_NUM_THREADS=1 BLIS_NUM_THREADS=1 VECLIB_MAXIMUM_THREADS=1 NUMBA_NUM_THREADS=1 NUMEXPR_NUM_THREADS=1 pytest flags: -m pytest tests --ignore-glob=tests/test_fwbw_overlap.py ============================= test session starts ============================== platform linux -- Python 3.13.15, pytest-9.1.1, pluggy-1.6.0 rootdir: /build/source configfile: pyproject.toml plugins: hydra-core-1.3.5, cov-stub-1.0.0, timeout-2.4.0, anyio-4.14.2 collected 3829 items tests/test_attention_patterns.py ....................................... [ 1%] ........................................................................ [ 2%] ......................... [ 3%] tests/test_checkpoint.py ................................sssssssssssssss [ 4%] sssssssssssssssss........ssssssssssssssssssssssssssssssssssssssss [ 6%] tests/test_fmha_flop_formula.py ssssssssssssssss. [ 6%] tests/test_fmha_merge_attentions.py ssssssssssssssssssssssssssssssssssss [ 7%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] sssssssssss [ 10%] tests/test_indexing.py sssssss [ 10%] tests/test_mem_eff_attention.py ssssssssssssssssssssssssssssssssssssssss [ 11%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 13%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 15%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 16%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 18%] ssssssssssssss.......sssssssssssssssssssssssssssssssssssssssssssssssssss [ 20%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 22%] ssssssssssssssssssssssssssssssssssssss.ssssss.sssssss..sssssssssssssssss [ 24%] ssssssssssssssssssssssssssssssssssssss.sssssssssssssssssssssssssssssssss [ 26%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 28%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 30%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 31%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 33%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 35%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 37%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 39%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 41%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 43%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 45%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 46%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 48%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 50%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 52%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 54%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 56%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 58%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 60%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 62%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 63%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 65%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 67%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 69%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 71%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 73%] ssssssssssssssss [ 73%] tests/test_multiprocessing_utils.py . [ 73%] tests/test_profiler.py ssssssssssssssss [ 74%] tests/test_rmsnorm.py ssssssssssssssssss [ 74%] tests/test_rope_padded.py ssssssssssssssssssssssssssssssssssssssssssssss [ 75%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 77%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 79%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 81%] sssssssssssssssssssssssssssss [ 82%] tests/test_seqpar.py sssssssssssssssssssssssssssssssssssssssssssssssssss [ 83%] sssssssssssssssssssssssssssssssssssssssssssss [ 84%] tests/test_sequence_parallel_fused_ops.py ssssssssssssssssssssssssssssss [ 85%] ssssssssssssssssssssssssssssssssssssssssssssss [ 86%] tests/test_sparse_tensors.py ..... [ 86%] tests/test_sparsity24.py sssssssssssssssssssssssssssssssssssssssssssssss [ 88%] sssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 89%] tests/test_splitk_reference.py .................. [ 90%] tests/test_tiled_matmul.py sssssssssssssssssssssssssssssssssssssssssssss [ 91%] sssssssssssssss [ 91%] tests/test_tree_attention.py sssssssssssssssssssssssssssssssssssssssssss [ 92%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 94%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 96%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 98%] sssssssssssss.................... [ 99%] tests/test_triton_varargs.py sssss [ 99%] tests/test_unbind.py .................... [100%] =============================== warnings summary =============================== tests/test_attention_patterns.py::test_local_1d_pattern[50-3] /nix/store/rgbk13h7ws87imm1mccvyvmpk34as3y4-python3.13-torch-2.13.0/lib/python3.13/site-packages/torch/functional.py:504: UserWarning: torch.meshgrid: in an upcoming release, it will be required to pass the indexing argument. (Triggered internally at /build/pytorch/aten/src/ATen/native/TensorShape.cpp:4215.) return _VF.meshgrid(tensors, **kwargs) # type: ignore[attr-defined] tests/test_checkpoint.py: 16 warnings /build/source/tests/test_checkpoint.py:67: FutureWarning: xformers.checkpoint.checkpoint is deprecated and will be removed in a future version of xFormers. Use torch.utils.checkpoint.checkpoint with use_reentrant=False and context_fn=functools.partial(torch.utils.checkpoint.create_selective_checkpoint_contexts, policy_fn). out = checkpoint(module[i], out, policy_fn=policy_fn) tests/test_checkpoint.py: 16 warnings /build/source/tests/test_checkpoint.py:100: FutureWarning: xformers.checkpoint.checkpoint is deprecated and will be removed in a future version of xFormers. Use torch.utils.checkpoint.checkpoint with use_reentrant=False and context_fn=functools.partial(torch.utils.checkpoint.create_selective_checkpoint_contexts, policy_fn). out = checkpoint(module, out, policy_fn=policy_fn) tests/test_checkpoint.py::test_list_operators /build/source/tests/test_checkpoint.py:195: FutureWarning: xformers.checkpoint.list_operators is deprecated and will be removed in a future version of xFormers. Inspect a region's operators with a custom torch.utils._python_dispatch.TorchDispatchMode. operators = list_operators(module, inputs) tests/test_tree_attention.py: 16 warnings /build/source/tests/test_tree_attention.py:458: FutureWarning: `torch.testing.assert_allclose()` is deprecated since 1.12 and will be removed in a future release. Please use `torch.testing.assert_close()` instead. You can find detailed upgrade instructions in https://github.com/pytorch/pytorch/issues/61844. torch.testing.assert_allclose( -- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html =============== 253 passed, 3576 skipped, 50 warnings in 15.04s ================ Finished executing pytestCheckPhase Running phase: pytestcachePhase Running phase: pytestRemoveBytecodePhase