python3.13-xformers-0.0.35
checks.aarch64-linux.machine-inference01
· build #7
· raw
1structuredAttrs is enabled2Sourcing python-remove-tests-dir-hook3Sourcing python-catch-conflicts-hook.sh4Sourcing python-remove-bin-bytecode-hook.sh5Sourcing pypa-build-hook6Using pypaBuildPhase7Sourcing python-runtime-deps-check-hook8Using pythonRuntimeDepsCheckHook9Sourcing pypa-install-hook10Using pypaInstallPhase11Sourcing python-imports-check-hook.sh12Using pythonImportsCheckPhase13Sourcing python-metadata-check-hook.sh14Using pythonMetadataCheckPhase15Sourcing python-namespaces-hook16Sourcing python-catch-conflicts-hook.sh17Sourcing setup-cuda-hook18Sourcing pytest-check-hook19Using pytestCheckPhase20Using checkPhaseThreadLimitHook21source: sourcing removeStubsFromRunpathHook.bash (hostOffset=0) (targetOffset=1)22source: added removeStubsFromRunpathHookRegistration to prePhases23Sourcing fix-elf-files.sh
removeStubsFromRunpathHookRegistration
24removeStubsFromRunpathHookRegistration: added removeStubsFromRunpath to postFixupHooksunpackPhase
25unpacking source archive /nix/store/wff5yymdwhnqc7hx5lcxp9hf5yvd6i89-source26source root is source27setting SOURCE_DATE_EPOCH to timestamp 315619200 of file "source/xformers/utils.py"patchPhase
28applying patch /nix/store/ai1qwfaq5vilgyl1sj9vxs1434ax99fp-xformers-fix-tests-pytorch-2.13.patch29patching file CHANGELOG.md30patching file xformers/checkpoint.pyconfigurePhase
31Executing setupCUDAToolkitCompilers32no configure script, doing nothingbuildPhase
33Executing pypaBuildPhase34Creating a wheel...35pypa build flags: --no-isolation --outdir dist/ --wheel36* Getting build dependencies for wheel...37/build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry38Error in cpuinfo: failed to parse the list of possible processors in /sys/devices/system/cpu/possible39Error in cpuinfo: failed to parse the list of present processors in /sys/devices/system/cpu/present40Error in cpuinfo: failed to parse both lists of possible and present processors41W0930 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'42/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.43 warn(44/nix/store/8p92j2ppflf7igibazxiwsh3wa8y528w-python3.13-setuptools-83.0.0/lib/python3.13/site-packages/setuptools/dist.py:765: SetuptoolsDeprecationWarning: License classifiers are deprecated.45!!4647 ********************************************************************************48 Please consider removing the following classifiers in favor of a SPDX license expression:4950 License :: OSI Approved :: BSD License5152 See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details.53 ********************************************************************************5455!!56 self._finalize_license_expression()57running egg_info58creating xformers.egg-info59writing xformers.egg-info/PKG-INFO60writing dependency_links to xformers.egg-info/dependency_links.txt61writing requirements to xformers.egg-info/requires.txt62writing top-level names to xformers.egg-info/top_level.txt63writing manifest file 'xformers.egg-info/SOURCES.txt'64W0930 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.65reading manifest file 'xformers.egg-info/SOURCES.txt'66reading manifest template 'MANIFEST.in'67adding license file 'LICENSE'68writing manifest file 'xformers.egg-info/SOURCES.txt'69* Building wheel...70/build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry71Error in cpuinfo: failed to parse the list of possible processors in /sys/devices/system/cpu/possible72Error in cpuinfo: failed to parse the list of present processors in /sys/devices/system/cpu/present73Error in cpuinfo: failed to parse both lists of possible and present processors74W0930 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'75/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.76 warn(77/nix/store/8p92j2ppflf7igibazxiwsh3wa8y528w-python3.13-setuptools-83.0.0/lib/python3.13/site-packages/setuptools/dist.py:765: SetuptoolsDeprecationWarning: License classifiers are deprecated.78!!7980 ********************************************************************************81 Please consider removing the following classifiers in favor of a SPDX license expression:8283 License :: OSI Approved :: BSD License8485 See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details.86 ********************************************************************************8788!!89 self._finalize_license_expression()90running bdist_wheel91W0930 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.92running build93running build_py94creating build/lib.linux-aarch64-cpython-313/xformers95copying xformers/test.py -> build/lib.linux-aarch64-cpython-313/xformers96copying xformers/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers97copying xformers/checkpoint.py -> build/lib.linux-aarch64-cpython-313/xformers98copying xformers/utils.py -> build/lib.linux-aarch64-cpython-313/xformers99copying xformers/attn_bias_utils.py -> build/lib.linux-aarch64-cpython-313/xformers100copying xformers/info.py -> build/lib.linux-aarch64-cpython-313/xformers101copying xformers/_deprecation_warning.py -> build/lib.linux-aarch64-cpython-313/xformers102copying xformers/_cpp_lib.py -> build/lib.linux-aarch64-cpython-313/xformers103copying xformers/fwbw_overlap.py -> build/lib.linux-aarch64-cpython-313/xformers104creating build/lib.linux-aarch64-cpython-313/xformers/ops105copying xformers/ops/sp24.py -> build/lib.linux-aarch64-cpython-313/xformers/ops106copying xformers/ops/rope_padded.py -> build/lib.linux-aarch64-cpython-313/xformers/ops107copying xformers/ops/rmsnorm.py -> build/lib.linux-aarch64-cpython-313/xformers/ops108copying xformers/ops/tiled_matmul.py -> build/lib.linux-aarch64-cpython-313/xformers/ops109copying xformers/ops/tree_attention.py -> build/lib.linux-aarch64-cpython-313/xformers/ops110copying xformers/ops/common.py -> build/lib.linux-aarch64-cpython-313/xformers/ops111copying xformers/ops/swiglu_op.py -> build/lib.linux-aarch64-cpython-313/xformers/ops112copying xformers/ops/seqpar.py -> build/lib.linux-aarch64-cpython-313/xformers/ops113copying xformers/ops/unbind.py -> build/lib.linux-aarch64-cpython-313/xformers/ops114copying xformers/ops/sequence_parallel_fused_ops.py -> build/lib.linux-aarch64-cpython-313/xformers/ops115copying xformers/ops/modpar_layers.py -> build/lib.linux-aarch64-cpython-313/xformers/ops116copying xformers/ops/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops117copying xformers/ops/indexing.py -> build/lib.linux-aarch64-cpython-313/xformers/ops118copying xformers/ops/differentiable_collectives.py -> build/lib.linux-aarch64-cpython-313/xformers/ops119creating build/lib.linux-aarch64-cpython-313/xformers/flash_attn_3120copying xformers/flash_attn_3/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/flash_attn_3121creating build/lib.linux-aarch64-cpython-313/xformers/components122copying xformers/components/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/components123creating build/lib.linux-aarch64-cpython-313/xformers/benchmarks124copying xformers/benchmarks/benchmark_mem_eff_attention.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks125copying xformers/benchmarks/benchmark_tiled_matmul.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks126copying xformers/benchmarks/benchmark_sp24.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks127copying xformers/benchmarks/benchmark_indexing.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks128copying xformers/benchmarks/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks129copying xformers/benchmarks/benchmark_merge_attentions.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks130copying xformers/benchmarks/benchmark_sequence_parallel_fused.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks131copying xformers/benchmarks/benchmark_attn_decoding.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks132copying xformers/benchmarks/utils.py -> build/lib.linux-aarch64-cpython-313/xformers/benchmarks133creating build/lib.linux-aarch64-cpython-313/xformers/sparse134copying xformers/sparse/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/sparse135copying xformers/sparse/blocksparse_tensor.py -> build/lib.linux-aarch64-cpython-313/xformers/sparse136copying xformers/sparse/utils.py -> build/lib.linux-aarch64-cpython-313/xformers/sparse137creating build/lib.linux-aarch64-cpython-313/xformers/profiler138copying xformers/profiler/profiler_dcgm_impl.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler139copying xformers/profiler/profiler.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler140copying xformers/profiler/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler141copying xformers/profiler/api.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler142copying xformers/profiler/profile_analyzer.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler143copying xformers/profiler/profiler_dcgm.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler144copying xformers/profiler/device_limits.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler145copying xformers/profiler/find_slowest.py -> build/lib.linux-aarch64-cpython-313/xformers/profiler146creating build/lib.linux-aarch64-cpython-313/xformers/triton147copying xformers/triton/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/triton148copying xformers/triton/vararg_kernel.py -> build/lib.linux-aarch64-cpython-313/xformers/triton149copying xformers/triton/importing.py -> build/lib.linux-aarch64-cpython-313/xformers/triton150creating build/lib.linux-aarch64-cpython-313/xformers/ops/fmha151copying xformers/ops/fmha/torch_attention_compat.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha152copying xformers/ops/fmha/cutlass_blackwell.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha153copying xformers/ops/fmha/attn_bias.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha154copying xformers/ops/fmha/dispatch.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha155copying xformers/ops/fmha/merge_training.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha156copying xformers/ops/fmha/flash.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha157copying xformers/ops/fmha/flash3.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha158copying xformers/ops/fmha/ck.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha159copying xformers/ops/fmha/triton_splitk.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha160copying xformers/ops/fmha/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha161copying xformers/ops/fmha/cutlass.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha162copying xformers/ops/fmha/ck_splitk.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha163copying xformers/ops/fmha/common.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha164creating build/lib.linux-aarch64-cpython-313/xformers/ops/_triton165copying xformers/ops/_triton/k_index_select_cat.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton166copying xformers/ops/_triton/rope_padded_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton167copying xformers/ops/_triton/rmsnorm_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton168copying xformers/ops/_triton/matmul_perf_model.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton169copying xformers/ops/_triton/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton170copying xformers/ops/_triton/tiled_matmul_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton171copying xformers/ops/_triton/k_scaled_index_add.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/_triton172creating build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton173copying xformers/ops/fmha/_triton/splitk_kernels.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton174copying xformers/ops/fmha/_triton/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton175creating build/lib.linux-aarch64-cpython-313/xformers/components/attention176copying xformers/components/attention/attention_patterns.py -> build/lib.linux-aarch64-cpython-313/xformers/components/attention177copying xformers/components/attention/__init__.py -> build/lib.linux-aarch64-cpython-313/xformers/components/attention178running build_ext179W0930 10:43:00.504000 117 torch/utils/cpp_extension.py:657] There are no g++ version bounds defined for CUDA version 13.2180building 'xformers._C' extension181creating build/temp.linux-aarch64-cpython-313/xformers/csrc/attention182creating build/temp.linux-aarch64-cpython-313/xformers/csrc/sparse24183g++ -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=_C184/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 gcc185ptxas info : 46 bytes gmem186ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'187ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv188 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads189ptxas info : Used 4 registers, used 0 barriers190ptxas info : Compile time = 2.433 ms191ptxas info : 46 bytes gmem192ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'193ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv194 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads195ptxas info : Used 4 registers, used 0 barriers196ptxas info : Compile time = 1.861 ms197/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 gcc198ptxas info : 46 bytes gmem199ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'200ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv201 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads202ptxas info : Used 4 registers, used 0 barriers203ptxas info : Compile time = 2.134 ms204ptxas info : 46 bytes gmem205ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'206ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv207 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads208ptxas info : Used 4 registers, used 0 barriers209ptxas info : Compile time = 1.854 ms210/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 gcc211ptxas info : 48 bytes gmem212ptxas 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'213ptxas 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_6ParamsE214 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads215ptxas info : Used 214 registers, used 1 barriers216ptxas info : Compile time = 171.988 ms217ptxas 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'218ptxas 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_6ParamsE219 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads220ptxas info : Used 220 registers, used 1 barriers221ptxas info : Compile time = 172.818 ms222ptxas 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'223ptxas 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_6ParamsE224 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads225ptxas info : Used 218 registers, used 1 barriers226ptxas info : Compile time = 164.534 ms227ptxas 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'228ptxas 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_6ParamsE229 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads230ptxas info : Used 224 registers, used 1 barriers231ptxas info : Compile time = 163.473 ms232ptxas 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'233ptxas 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_6ParamsE234 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads235ptxas info : Used 214 registers, used 1 barriers236ptxas info : Compile time = 170.488 ms237ptxas 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'238ptxas 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_6ParamsE239 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads240ptxas info : Used 220 registers, used 1 barriers241ptxas info : Compile time = 174.017 ms242ptxas 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'243ptxas 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_6ParamsE244 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads245ptxas info : Used 218 registers, used 1 barriers246ptxas info : Compile time = 169.267 ms247ptxas 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'248ptxas 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_6ParamsE249 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads250ptxas info : Used 226 registers, used 1 barriers251ptxas info : Compile time = 166.935 ms252ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'253ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv254 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads255ptxas info : Used 4 registers, used 0 barriers256ptxas info : Compile time = 0.700 ms257ptxas info : 48 bytes gmem258ptxas 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'259ptxas 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_6ParamsE260 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads261ptxas info : Used 214 registers, used 1 barriers262ptxas info : Compile time = 172.816 ms263ptxas 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'264ptxas 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_6ParamsE265 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads266ptxas info : Used 220 registers, used 1 barriers267ptxas info : Compile time = 173.837 ms268ptxas 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'269ptxas 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_6ParamsE270 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads271ptxas info : Used 218 registers, used 1 barriers272ptxas info : Compile time = 165.266 ms273ptxas 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'274ptxas 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_6ParamsE275 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads276ptxas info : Used 224 registers, used 1 barriers277ptxas info : Compile time = 163.301 ms278ptxas 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'279ptxas 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_6ParamsE280 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads281ptxas info : Used 214 registers, used 1 barriers282ptxas info : Compile time = 171.595 ms283ptxas 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'284ptxas 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_6ParamsE285 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads286ptxas info : Used 220 registers, used 1 barriers287ptxas info : Compile time = 175.636 ms288ptxas 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'289ptxas 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_6ParamsE290 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads291ptxas info : Used 218 registers, used 1 barriers292ptxas info : Compile time = 169.881 ms293ptxas 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'294ptxas 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_6ParamsE295 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads296ptxas info : Used 226 registers, used 1 barriers297ptxas info : Compile time = 168.442 ms298ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'299ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv300 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads301ptxas info : Used 4 registers, used 0 barriers302ptxas info : Compile time = 0.771 ms303/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 gcc304ptxas info : 48 bytes gmem305ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'306ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv307 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads308ptxas info : Used 4 registers, used 0 barriers309ptxas info : Compile time = 2.116 ms310ptxas info : 48 bytes gmem311ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'312ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv313 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads314ptxas info : Used 4 registers, used 0 barriers315ptxas info : Compile time = 2.793 ms316g++ -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=_C317/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 gcc318ptxas info : 46 bytes gmem319ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE' for 'sm_121'320ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE321 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads322ptxas info : Used 96 registers, used 0 barriers323ptxas info : Compile time = 52.814 ms324ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE' for 'sm_121'325ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE326 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads327ptxas info : Used 96 registers, used 0 barriers328ptxas info : Compile time = 50.078 ms329ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'330ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv331 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads332ptxas info : Used 4 registers, used 0 barriers333ptxas info : Compile time = 0.588 ms334ptxas info : 46 bytes gmem335ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE' for 'sm_120'336ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEEEEvNT_6ParamsE337 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads338ptxas info : Used 96 registers, used 0 barriers339ptxas info : Compile time = 54.318 ms340ptxas info : Compiling entry function '_ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE' for 'sm_120'341ptxas info : Function properties for _ZN54_GLOBAL__N__5361b14d_17_sparse24_apply_cu_4cca563f_23421sparse24_apply_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEEEEvNT_6ParamsE342 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads343ptxas info : Used 96 registers, used 0 barriers344ptxas info : Compile time = 49.700 ms345ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'346ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv347 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads348ptxas info : Used 4 registers, used 0 barriers349ptxas info : Compile time = 0.577 ms350/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 gcc351ptxas info : 46 bytes gmem352ptxas 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'353ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb0ELb1EEEvNS_6ParamsIT_EE354 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads355ptxas info : Used 80 registers, used 0 barriers356ptxas info : Compile time = 39.943 ms357ptxas 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'358ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb1ELb1EEEvNS_6ParamsIT_EE359 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads360ptxas info : Used 59 registers, used 0 barriers361ptxas info : Compile time = 45.983 ms362ptxas 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'363ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb0ELb1EEEvNS_6ParamsIT_EE364 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads365ptxas info : Used 80 registers, used 0 barriers366ptxas info : Compile time = 38.088 ms367ptxas 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'368ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb1ELb1EEEvNS_6ParamsIT_EE369 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads370ptxas info : Used 59 registers, used 0 barriers371ptxas info : Compile time = 46.816 ms372ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'373ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv374 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads375ptxas info : Used 4 registers, used 0 barriers376ptxas info : Compile time = 0.636 ms377ptxas info : 46 bytes gmem378ptxas 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'379ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb0ELb1EEEvNS_6ParamsIT_EE380 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads381ptxas info : Used 80 registers, used 0 barriers382ptxas info : Compile time = 39.949 ms383ptxas 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'384ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass10bfloat16_tELb1ELb1EEEvNS_6ParamsIT_EE385 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads386ptxas info : Used 59 registers, used 0 barriers387ptxas info : Compile time = 45.869 ms388ptxas 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'389ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb0ELb1EEEvNS_6ParamsIT_EE390 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads391ptxas info : Used 80 registers, used 0 barriers392ptxas info : Compile time = 37.928 ms393ptxas 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'394ptxas info : Function properties for _ZN67_GLOBAL__N__666b03f5_30_sparse24_apply_dense_output_cu_4cca563f_25729sparse24_apply_dense_output_kIN7cutlass6half_tELb1ELb1EEEvNS_6ParamsIT_EE395 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads396ptxas info : Used 59 registers, used 0 barriers397ptxas info : Compile time = 46.791 ms398ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'399ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv400 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads401ptxas info : Used 4 registers, used 0 barriers402ptxas info : Compile time = 0.620 ms403/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 gcc404xformers/csrc/sparse24/sparse24_gemm_sm90.cu(139): warning #177-D: member "<unnamed>::SparseRowwiseKernel<cutlass::bfloat16_t>::kElementOutAt" was declared but never referenced405 static constexpr auto kElementOutAt = torch::headeronly::ScalarType::BFloat16;406 ^407408Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"409410xformers/csrc/sparse24/sparse24_gemm_sm90.cu(139): warning #177-D: member "<unnamed>::SparseRowwiseKernel<cutlass::bfloat16_t>::kElementOutAt" was declared but never referenced411 static constexpr auto kElementOutAt = torch::headeronly::ScalarType::BFloat16;412 ^413414Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"415416ptxas info : 155 bytes gmem417ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'418ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv419 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads420ptxas info : Used 4 registers, used 0 barriers421ptxas info : Compile time = 1.821 ms422ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE' for 'sm_121'423ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE424 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads425ptxas info : Used 96 registers, used 1 barriers426ptxas info : Compile time = 109.399 ms427ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE' for 'sm_121'428ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE429 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads430ptxas info : Used 162 registers, used 1 barriers431ptxas info : Compile time = 240.021 ms432ptxas 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'433ptxas 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_6ParamsE434 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads435ptxas info : Used 24 registers, used 0 barriers436ptxas info : Compile time = 1.500 ms437ptxas info : 155 bytes gmem438ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'439ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv440 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads441ptxas info : Used 4 registers, used 0 barriers442ptxas info : Compile time = 1.639 ms443ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE' for 'sm_120'444ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_10bfloat16_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi64EEEEEEEEEvNT_6ParamsE445 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads446ptxas info : Used 96 registers, used 1 barriers447ptxas info : Compile time = 109.581 ms448ptxas info : Compiling entry function '_ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE' for 'sm_120'449ptxas info : Function properties for _ZN7cutlass13device_kernelINS_9transform6kernel30SM90StructuredSparseCompressorIN4cute5tupleIJiiiiEEENS_12float_e4m3_tENS_6layout8RowMajorENS_20Sm90GemmSparseConfigINS4_11sparse_elemILi2ES7_EELNS4_4SM904GMMA5MajorE0ENSB_ILi8EhEENS4_1CILi128EEEEEEEEEvNT_6ParamsE450 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads451ptxas info : Used 162 registers, used 1 barriers452ptxas info : Compile time = 240.732 ms453ptxas 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'454ptxas 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_6ParamsE455 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads456ptxas info : Used 24 registers, used 0 barriers457ptxas info : Compile time = 1.556 ms458/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 gcc459ptxas info : 46 bytes gmem460ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'461ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv462 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads463ptxas info : Used 4 registers, used 0 barriers464ptxas info : Compile time = 1.753 ms465ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE' for 'sm_120'466ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE467 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads468ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size469ptxas info : Compile time = 134.253 ms470ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE' for 'sm_120'471ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE472 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads473ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size474ptxas info : Compile time = 136.809 ms475ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE' for 'sm_120'476ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE477 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads478ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size479ptxas info : Compile time = 51.122 ms480ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE' for 'sm_120'481ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE482 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads483ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size484ptxas info : Compile time = 54.085 ms485ptxas info : 46 bytes gmem486ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'487ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv488 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads489ptxas info : Used 4 registers, used 0 barriers490ptxas info : Compile time = 1.708 ms491ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE' for 'sm_121'492ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb1EEEvNT_6ParamsE493 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads494ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size495ptxas info : Compile time = 135.907 ms496ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE' for 'sm_121'497ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb1EEEvNT_6ParamsE498 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads499ptxas info : Used 68 registers, used 0 barriers, 8 bytes cumulative stack size500ptxas info : Compile time = 138.598 ms501ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE' for 'sm_121'502ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass10bfloat16_tELb0EEEvNT_6ParamsE503 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads504ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size505ptxas info : Compile time = 51.033 ms506ptxas info : Compiling entry function '_Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE' for 'sm_121'507ptxas info : Function properties for _Z27sparse24_largest_mask_2d_cuI14Sp24MaskKernelIN7cutlass6half_tELb0EEEvNT_6ParamsE508 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads509ptxas info : Used 56 registers, used 0 barriers, 8 bytes cumulative stack size510ptxas info : Compile time = 54.259 ms511/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 gcc512ptxas info : 46 bytes gmem513ptxas 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'514ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_515 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads516ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size517ptxas info : Compile time = 314.039 ms518ptxas 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'519ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_520 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads521ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size522ptxas info : Compile time = 126.251 ms523ptxas 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'524ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_525 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads526ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size527ptxas info : Compile time = 126.670 ms528ptxas 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'529ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_530 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads531ptxas info : Used 96 registers, used 0 barriers532ptxas info : Compile time = 156.318 ms533ptxas 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'534ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_535 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads536ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size537ptxas info : Compile time = 293.931 ms538ptxas 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'539ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_540 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads541ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size542ptxas info : Compile time = 309.011 ms543ptxas 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'544ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_545 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads546ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size547ptxas info : Compile time = 126.945 ms548ptxas 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'549ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_550 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads551ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size552ptxas info : Compile time = 126.869 ms553ptxas 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'554ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_555 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads556ptxas info : Used 96 registers, used 0 barriers557ptxas info : Compile time = 154.925 ms558ptxas 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'559ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_560 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads561ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size562ptxas info : Compile time = 304.721 ms563ptxas 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'564ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_565 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads566ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size567ptxas info : Compile time = 310.663 ms568ptxas 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'569ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_570 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads571ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size572ptxas info : Compile time = 124.588 ms573ptxas 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'574ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_575 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads576ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size577ptxas info : Compile time = 125.155 ms578ptxas 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'579ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_580 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads581ptxas info : Used 96 registers, used 0 barriers582ptxas info : Compile time = 153.806 ms583ptxas 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'584ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_585 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads586ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size587ptxas info : Compile time = 292.594 ms588ptxas 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'589ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_590 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads591ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size592ptxas info : Compile time = 306.241 ms593ptxas 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'594ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_595 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads596ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size597ptxas info : Compile time = 126.673 ms598ptxas 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'599ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_600 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads601ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size602ptxas info : Compile time = 126.637 ms603ptxas 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'604ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_605 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads606ptxas info : Used 96 registers, used 0 barriers607ptxas info : Compile time = 154.419 ms608ptxas 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'609ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_610 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads611ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size612ptxas info : Compile time = 301.579 ms613ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'614ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv615 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads616ptxas info : Used 4 registers, used 0 barriers617ptxas info : Compile time = 0.784 ms618ptxas info : 46 bytes gmem619ptxas 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'620ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_621 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads622ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size623ptxas info : Compile time = 315.079 ms624ptxas 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'625ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_626 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads627ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size628ptxas info : Compile time = 126.593 ms629ptxas 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'630ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_631 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads632ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size633ptxas info : Compile time = 127.147 ms634ptxas 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'635ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_636 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads637ptxas info : Used 96 registers, used 0 barriers638ptxas info : Compile time = 156.576 ms639ptxas 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'640ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_641 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads642ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size643ptxas info : Compile time = 294.921 ms644ptxas 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'645ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_646 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads647ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size648ptxas info : Compile time = 309.906 ms649ptxas 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'650ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_651 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads652ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size653ptxas info : Compile time = 126.965 ms654ptxas 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'655ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_656 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads657ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size658ptxas info : Compile time = 126.687 ms659ptxas 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'660ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_661 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads662ptxas info : Used 96 registers, used 0 barriers663ptxas info : Compile time = 155.787 ms664ptxas 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'665ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass10bfloat16_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_666 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads667ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size668ptxas info : Compile time = 306.276 ms669ptxas 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'670ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_671 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads672ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size673ptxas info : Compile time = 310.619 ms674ptxas 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'675ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_676 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads677ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size678ptxas info : Compile time = 124.817 ms679ptxas 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'680ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_681 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads682ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size683ptxas info : Compile time = 125.454 ms684ptxas 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'685ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_686 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads687ptxas info : Used 96 registers, used 0 barriers688ptxas info : Compile time = 154.234 ms689ptxas 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'690ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_19MetadataCutlassSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_691 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads692ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size693ptxas info : Compile time = 292.750 ms694ptxas 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'695ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_5AbsOpEEEEEvNT_6ParamsET0_T1_696 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads697ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size698ptxas info : Compile time = 306.190 ms699ptxas 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'700ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_5AbsOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_701 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads702ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size703ptxas info : Compile time = 126.516 ms704ptxas 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'705ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi2ELi2ELi2ELi2ELb0EEEEEvNT_6ParamsET0_T1_706 24 bytes stack frame, 20 bytes spill stores, 20 bytes spill loads707ptxas info : Used 96 registers, used 0 barriers, 24 bytes cumulative stack size708ptxas info : Compile time = 127.104 ms709ptxas 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'710ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_10LineByLineINS2_10IdentityOpELi1ELi1ELi2ELi2ELb1EEEEEvNT_6ParamsET0_T1_711 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads712ptxas info : Used 96 registers, used 0 barriers713ptxas info : Compile time = 154.725 ms714ptxas 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'715ptxas info : Function properties for _ZN53_GLOBAL__N__b60c2222_16_sparse24_pack_cu_4cca563f_32634sparse24_sparsify_both_ways_kernelIN8xformers4sp2411KernelTypesIN7cutlass6half_tEEENS2_22MetadataCuSparseLtSm80ENS2_19LargestValuesGreedyINS2_10IdentityOpEEEEEvNT_6ParamsET0_T1_716 16 bytes stack frame, 16 bytes spill stores, 16 bytes spill loads717ptxas info : Used 96 registers, used 0 barriers, 16 bytes cumulative stack size718ptxas info : Compile time = 303.476 ms719ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'720ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv721 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads722ptxas info : Used 4 registers, used 0 barriers723ptxas info : Compile time = 0.791 ms724/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 gcc725ptxas info : 46 bytes gmem726ptxas 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'727ptxas info : Function properties for _ZN58_GLOBAL__N__0b8b56aa_21_sparse24_pack_test_cu_4cca563f_34924meta_shuffle_test_kernelEN5torch10headeronly6detail27GenericPackedTensorAccessorINS2_14TensorAccessorIN3c1018HeaderOnlyArrayRefIlEElLm2ENS1_16DefaultPtrTraitsElEENS2_26HeaderOnlyIndexBoundsCheckILm3ElEElLm3ES8_lEESC_b728 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads729ptxas info : Used 26 registers, used 0 barriers730ptxas info : Compile time = 10.000 ms731ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'732ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv733 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads734ptxas info : Used 4 registers, used 0 barriers735ptxas info : Compile time = 0.886 ms736ptxas info : 46 bytes gmem737ptxas 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'738ptxas info : Function properties for _ZN58_GLOBAL__N__0b8b56aa_21_sparse24_pack_test_cu_4cca563f_34924meta_shuffle_test_kernelEN5torch10headeronly6detail27GenericPackedTensorAccessorINS2_14TensorAccessorIN3c1018HeaderOnlyArrayRefIlEElLm2ENS1_16DefaultPtrTraitsElEENS2_26HeaderOnlyIndexBoundsCheckILm3ElEElLm3ES8_lEESC_b739 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads740ptxas info : Used 26 registers, used 0 barriers741ptxas info : Compile time = 9.690 ms742ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'743ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv744 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads745ptxas info : Used 4 registers, used 0 barriers746ptxas info : Compile time = 0.894 ms747/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 gcc748ptxas info : 46 bytes gmem749ptxas 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'750ptxas 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_751 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads752ptxas info : Used 24 registers, used 0 barriers753ptxas info : Compile time = 84.550 ms754ptxas 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'755ptxas 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_756 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads757ptxas info : Used 23 registers, used 0 barriers758ptxas info : Compile time = 71.719 ms759ptxas 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'760ptxas 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_761 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads762ptxas info : Used 32 registers, used 0 barriers763ptxas info : Compile time = 87.689 ms764ptxas 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'765ptxas 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_766 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads767ptxas info : Used 28 registers, used 0 barriers768ptxas info : Compile time = 77.425 ms769ptxas 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'770ptxas 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_771 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads772ptxas info : Used 24 registers, used 0 barriers773ptxas info : Compile time = 81.265 ms774ptxas 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'775ptxas 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_776 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads777ptxas info : Used 23 registers, used 0 barriers778ptxas info : Compile time = 68.931 ms779ptxas 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'780ptxas 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_781 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads782ptxas info : Used 32 registers, used 0 barriers783ptxas info : Compile time = 85.886 ms784ptxas 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'785ptxas 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_786 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads787ptxas info : Used 28 registers, used 0 barriers788ptxas info : Compile time = 73.455 ms789ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_120'790ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv791 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads792ptxas info : Used 4 registers, used 0 barriers793ptxas info : Compile time = 0.676 ms794ptxas info : 46 bytes gmem795ptxas 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'796ptxas 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_797 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads798ptxas info : Used 24 registers, used 0 barriers799ptxas info : Compile time = 84.736 ms800ptxas 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'801ptxas 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_802 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads803ptxas info : Used 23 registers, used 0 barriers804ptxas info : Compile time = 71.683 ms805ptxas 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'806ptxas 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_807 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads808ptxas info : Used 32 registers, used 0 barriers809ptxas info : Compile time = 87.624 ms810ptxas 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'811ptxas 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_812 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads813ptxas info : Used 28 registers, used 0 barriers814ptxas info : Compile time = 77.222 ms815ptxas 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'816ptxas 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_817 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads818ptxas info : Used 24 registers, used 0 barriers819ptxas info : Compile time = 81.803 ms820ptxas 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'821ptxas 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_822 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads823ptxas info : Used 23 registers, used 0 barriers824ptxas info : Compile time = 68.885 ms825ptxas 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'826ptxas 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_827 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads828ptxas info : Used 32 registers, used 0 barriers829ptxas info : Compile time = 85.740 ms830ptxas 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'831ptxas 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_832 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads833ptxas info : Used 28 registers, used 0 barriers834ptxas info : Compile time = 73.030 ms835ptxas info : Compiling entry function '_ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv' for 'sm_121'836ptxas info : Function properties for _ZN3cub22_V_300200_SM_1200_12106detail11EmptyKernelIvEEvv837 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads838ptxas info : Used 4 registers, used 0 barriers839ptxas info : Compile time = 0.688 ms840g++ -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.so841installing to build/bdist.linux-aarch64/wheel842running install843running install_lib844creating build/bdist.linux-aarch64/wheel845creating build/bdist.linux-aarch64/wheel/xformers846copying build/lib.linux-aarch64-cpython-313/xformers/test.py -> build/bdist.linux-aarch64/wheel/./xformers847copying build/lib.linux-aarch64-cpython-313/xformers/_deprecation_warning.py -> build/bdist.linux-aarch64/wheel/./xformers848creating build/bdist.linux-aarch64/wheel/xformers/benchmarks849copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_sp24.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks850copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_merge_attentions.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks851copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_attn_decoding.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks852copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_mem_eff_attention.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks853copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_tiled_matmul.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks854copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks855copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/utils.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks856copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_indexing.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks857copying build/lib.linux-aarch64-cpython-313/xformers/benchmarks/benchmark_sequence_parallel_fused.py -> build/bdist.linux-aarch64/wheel/./xformers/benchmarks858copying build/lib.linux-aarch64-cpython-313/xformers/utils.py -> build/bdist.linux-aarch64/wheel/./xformers859copying build/lib.linux-aarch64-cpython-313/xformers/checkpoint.py -> build/bdist.linux-aarch64/wheel/./xformers860copying build/lib.linux-aarch64-cpython-313/xformers/fwbw_overlap.py -> build/bdist.linux-aarch64/wheel/./xformers861copying build/lib.linux-aarch64-cpython-313/xformers/_C.so -> build/bdist.linux-aarch64/wheel/./xformers862copying build/lib.linux-aarch64-cpython-313/xformers/_cpp_lib.py -> build/bdist.linux-aarch64/wheel/./xformers863creating build/bdist.linux-aarch64/wheel/xformers/ops864copying build/lib.linux-aarch64-cpython-313/xformers/ops/rmsnorm.py -> build/bdist.linux-aarch64/wheel/./xformers/ops865copying build/lib.linux-aarch64-cpython-313/xformers/ops/swiglu_op.py -> build/bdist.linux-aarch64/wheel/./xformers/ops866creating build/bdist.linux-aarch64/wheel/xformers/ops/fmha867copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/cutlass.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha868copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/triton_splitk.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha869creating build/bdist.linux-aarch64/wheel/xformers/ops/fmha/_triton870copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha/_triton871copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/_triton/splitk_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha/_triton872copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/ck_splitk.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha873copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha874copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/ck.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha875copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/dispatch.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha876copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/attn_bias.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha877copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/flash3.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha878copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/flash.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha879copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/cutlass_blackwell.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha880copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/merge_training.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha881copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/common.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha882copying build/lib.linux-aarch64-cpython-313/xformers/ops/fmha/torch_attention_compat.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/fmha883copying build/lib.linux-aarch64-cpython-313/xformers/ops/common.py -> build/bdist.linux-aarch64/wheel/./xformers/ops884copying build/lib.linux-aarch64-cpython-313/xformers/ops/modpar_layers.py -> build/bdist.linux-aarch64/wheel/./xformers/ops885copying build/lib.linux-aarch64-cpython-313/xformers/ops/tree_attention.py -> build/bdist.linux-aarch64/wheel/./xformers/ops886copying build/lib.linux-aarch64-cpython-313/xformers/ops/indexing.py -> build/bdist.linux-aarch64/wheel/./xformers/ops887copying build/lib.linux-aarch64-cpython-313/xformers/ops/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops888copying build/lib.linux-aarch64-cpython-313/xformers/ops/unbind.py -> build/bdist.linux-aarch64/wheel/./xformers/ops889copying build/lib.linux-aarch64-cpython-313/xformers/ops/differentiable_collectives.py -> build/bdist.linux-aarch64/wheel/./xformers/ops890copying build/lib.linux-aarch64-cpython-313/xformers/ops/tiled_matmul.py -> build/bdist.linux-aarch64/wheel/./xformers/ops891copying build/lib.linux-aarch64-cpython-313/xformers/ops/sp24.py -> build/bdist.linux-aarch64/wheel/./xformers/ops892copying build/lib.linux-aarch64-cpython-313/xformers/ops/seqpar.py -> build/bdist.linux-aarch64/wheel/./xformers/ops893copying build/lib.linux-aarch64-cpython-313/xformers/ops/rope_padded.py -> build/bdist.linux-aarch64/wheel/./xformers/ops894copying build/lib.linux-aarch64-cpython-313/xformers/ops/sequence_parallel_fused_ops.py -> build/bdist.linux-aarch64/wheel/./xformers/ops895creating build/bdist.linux-aarch64/wheel/xformers/ops/_triton896copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/k_scaled_index_add.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton897copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/matmul_perf_model.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton898copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/rope_padded_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton899copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton900copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/rmsnorm_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton901copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/tiled_matmul_kernels.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton902copying build/lib.linux-aarch64-cpython-313/xformers/ops/_triton/k_index_select_cat.py -> build/bdist.linux-aarch64/wheel/./xformers/ops/_triton903creating build/bdist.linux-aarch64/wheel/xformers/flash_attn_3904copying build/lib.linux-aarch64-cpython-313/xformers/flash_attn_3/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/flash_attn_3905copying build/lib.linux-aarch64-cpython-313/xformers/info.py -> build/bdist.linux-aarch64/wheel/./xformers906copying build/lib.linux-aarch64-cpython-313/xformers/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers907copying build/lib.linux-aarch64-cpython-313/xformers/cpp_lib.json -> build/bdist.linux-aarch64/wheel/./xformers908creating build/bdist.linux-aarch64/wheel/xformers/triton909copying build/lib.linux-aarch64-cpython-313/xformers/triton/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/triton910copying build/lib.linux-aarch64-cpython-313/xformers/triton/importing.py -> build/bdist.linux-aarch64/wheel/./xformers/triton911copying build/lib.linux-aarch64-cpython-313/xformers/triton/vararg_kernel.py -> build/bdist.linux-aarch64/wheel/./xformers/triton912copying build/lib.linux-aarch64-cpython-313/xformers/attn_bias_utils.py -> build/bdist.linux-aarch64/wheel/./xformers913copying build/lib.linux-aarch64-cpython-313/xformers/version.py -> build/bdist.linux-aarch64/wheel/./xformers914creating build/bdist.linux-aarch64/wheel/xformers/profiler915copying build/lib.linux-aarch64-cpython-313/xformers/profiler/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler916copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profiler_dcgm_impl.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler917copying build/lib.linux-aarch64-cpython-313/xformers/profiler/find_slowest.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler918copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profiler.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler919copying build/lib.linux-aarch64-cpython-313/xformers/profiler/api.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler920copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profile_analyzer.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler921copying build/lib.linux-aarch64-cpython-313/xformers/profiler/profiler_dcgm.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler922copying build/lib.linux-aarch64-cpython-313/xformers/profiler/device_limits.py -> build/bdist.linux-aarch64/wheel/./xformers/profiler923creating build/bdist.linux-aarch64/wheel/xformers/components924creating build/bdist.linux-aarch64/wheel/xformers/components/attention925copying build/lib.linux-aarch64-cpython-313/xformers/components/attention/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/components/attention926copying build/lib.linux-aarch64-cpython-313/xformers/components/attention/attention_patterns.py -> build/bdist.linux-aarch64/wheel/./xformers/components/attention927copying build/lib.linux-aarch64-cpython-313/xformers/components/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/components928creating build/bdist.linux-aarch64/wheel/xformers/sparse929copying build/lib.linux-aarch64-cpython-313/xformers/sparse/utils.py -> build/bdist.linux-aarch64/wheel/./xformers/sparse930copying build/lib.linux-aarch64-cpython-313/xformers/sparse/__init__.py -> build/bdist.linux-aarch64/wheel/./xformers/sparse931copying build/lib.linux-aarch64-cpython-313/xformers/sparse/blocksparse_tensor.py -> build/bdist.linux-aarch64/wheel/./xformers/sparse932running install_egg_info933running egg_info934writing xformers.egg-info/PKG-INFO935writing dependency_links to xformers.egg-info/dependency_links.txt936writing requirements to xformers.egg-info/requires.txt937writing top-level names to xformers.egg-info/top_level.txt938reading manifest file 'xformers.egg-info/SOURCES.txt'939reading manifest template 'MANIFEST.in'940adding license file 'LICENSE'941writing manifest file 'xformers.egg-info/SOURCES.txt'942Copying xformers.egg-info to build/bdist.linux-aarch64/wheel/./xformers-0.0.35-py3.13.egg-info943running install_scripts944creating build/bdist.linux-aarch64/wheel/xformers-0.0.35.dist-info/WHEEL945creating '/build/source/dist/.tmp-8e8qf1el/xformers-0.0.35-py39-none-linux_aarch64.whl' and adding 'build/bdist.linux-aarch64/wheel' to it946adding 'xformers/_C.so'947adding 'xformers/__init__.py'948adding 'xformers/_cpp_lib.py'949adding 'xformers/_deprecation_warning.py'950adding 'xformers/attn_bias_utils.py'951adding 'xformers/checkpoint.py'952adding 'xformers/cpp_lib.json'953adding 'xformers/fwbw_overlap.py'954adding 'xformers/info.py'955adding 'xformers/test.py'956adding 'xformers/utils.py'957adding 'xformers/version.py'958adding 'xformers/benchmarks/__init__.py'959adding 'xformers/benchmarks/benchmark_attn_decoding.py'960adding 'xformers/benchmarks/benchmark_indexing.py'961adding 'xformers/benchmarks/benchmark_mem_eff_attention.py'962adding 'xformers/benchmarks/benchmark_merge_attentions.py'963adding 'xformers/benchmarks/benchmark_sequence_parallel_fused.py'964adding 'xformers/benchmarks/benchmark_sp24.py'965adding 'xformers/benchmarks/benchmark_tiled_matmul.py'966adding 'xformers/benchmarks/utils.py'967adding 'xformers/components/__init__.py'968adding 'xformers/components/attention/__init__.py'969adding 'xformers/components/attention/attention_patterns.py'970adding 'xformers/flash_attn_3/__init__.py'971adding 'xformers/ops/__init__.py'972adding 'xformers/ops/common.py'973adding 'xformers/ops/differentiable_collectives.py'974adding 'xformers/ops/indexing.py'975adding 'xformers/ops/modpar_layers.py'976adding 'xformers/ops/rmsnorm.py'977adding 'xformers/ops/rope_padded.py'978adding 'xformers/ops/seqpar.py'979adding 'xformers/ops/sequence_parallel_fused_ops.py'980adding 'xformers/ops/sp24.py'981adding 'xformers/ops/swiglu_op.py'982adding 'xformers/ops/tiled_matmul.py'983adding 'xformers/ops/tree_attention.py'984adding 'xformers/ops/unbind.py'985adding 'xformers/ops/_triton/__init__.py'986adding 'xformers/ops/_triton/k_index_select_cat.py'987adding 'xformers/ops/_triton/k_scaled_index_add.py'988adding 'xformers/ops/_triton/matmul_perf_model.py'989adding 'xformers/ops/_triton/rmsnorm_kernels.py'990adding 'xformers/ops/_triton/rope_padded_kernels.py'991adding 'xformers/ops/_triton/tiled_matmul_kernels.py'992adding 'xformers/ops/fmha/__init__.py'993adding 'xformers/ops/fmha/attn_bias.py'994adding 'xformers/ops/fmha/ck.py'995adding 'xformers/ops/fmha/ck_splitk.py'996adding 'xformers/ops/fmha/common.py'997adding 'xformers/ops/fmha/cutlass.py'998adding 'xformers/ops/fmha/cutlass_blackwell.py'999adding 'xformers/ops/fmha/dispatch.py'1000adding 'xformers/ops/fmha/flash.py'1001adding 'xformers/ops/fmha/flash3.py'1002adding 'xformers/ops/fmha/merge_training.py'1003adding 'xformers/ops/fmha/torch_attention_compat.py'1004adding 'xformers/ops/fmha/triton_splitk.py'1005adding 'xformers/ops/fmha/_triton/__init__.py'1006adding 'xformers/ops/fmha/_triton/splitk_kernels.py'1007adding 'xformers/profiler/__init__.py'1008adding 'xformers/profiler/api.py'1009adding 'xformers/profiler/device_limits.py'1010adding 'xformers/profiler/find_slowest.py'1011adding 'xformers/profiler/profile_analyzer.py'1012adding 'xformers/profiler/profiler.py'1013adding 'xformers/profiler/profiler_dcgm.py'1014adding 'xformers/profiler/profiler_dcgm_impl.py'1015adding 'xformers/sparse/__init__.py'1016adding 'xformers/sparse/blocksparse_tensor.py'1017adding 'xformers/sparse/utils.py'1018adding 'xformers/triton/__init__.py'1019adding 'xformers/triton/importing.py'1020adding 'xformers/triton/vararg_kernel.py'1021adding 'xformers-0.0.35.dist-info/licenses/LICENSE'1022adding 'xformers-0.0.35.dist-info/METADATA'1023adding 'xformers-0.0.35.dist-info/WHEEL'1024adding 'xformers-0.0.35.dist-info/top_level.txt'1025adding 'xformers-0.0.35.dist-info/RECORD'1026removing build/bdist.linux-aarch64/wheel1027Successfully built xformers-0.0.35-py39-none-linux_aarch64.whl1028Finished creating a wheel...1029Finished executing pypaBuildPhase1030buildPhase completed in 3 minutes 13 secondspythonRuntimeDepsCheckHook
1031Executing pythonRuntimeDepsCheck1032Checking runtime dependencies for xformers-0.0.35-py39-none-linux_aarch64.whl1033Finished executing pythonRuntimeDepsCheckinstallPhase
1034Executing pypaInstallPhase1035Successfully installed xformers-0.0.35-py39-none-linux_aarch64.whl1036Finished executing pypaInstallPhasepythonOutputDistPhase
1037Executing pythonOutputDistPhase1038Finished executing pythonOutputDistPhasefixupPhase
1039shrinking RPATHs of ELF executables and libraries in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist1040checking for references to /build/ in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist...1041patching script interpreter paths in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist1042shrinking RPATHs of ELF executables and libraries in /nix/store/82i0syl3b1d4sk9kx082xa9fdbgivphd-python3.13-xformers-0.0.35-dist1043shrinking RPATHs of ELF executables and libraries in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.351044shrinking /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib/python3.13/site-packages/xformers/_C.so1045checking for references to /build/ in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35...1046patching script interpreter paths in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.351047stripping (with command strip and flags -S -p) in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib1048shrinking RPATHs of ELF executables and libraries in /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.351049shrinking /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib/python3.13/site-packages/xformers/_C.so1050Executing pythonRemoveTestsDir1051Finished executing pythonRemoveTestsDir1052autoFixElfFiles: using removeStubsFromRunpath to fix /nix/store/4wpf3wkzv97y88ady453srv3vjsxfb7d-python3.13-xformers-0.0.35/lib/python3.13/site-packages/xformers/_C.soinstallCheckPhase
1053setting known thread variables to default: 11054 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=11055no Makefile or custom installCheckPhase, doing nothingpythonImportsCheckPhase
1056Executing pythonImportsCheckPhase1057Check whether the following modules can be imported: xformers xformers.components xformers.components.attention1058/build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entrypythonMetadataCheckPhase
1059Executing pythonMetadataCheckPhasepytestCheckPhase
1060Executing pytestCheckPhase1061/build/pytorch/third_party/ideep/mkl-dnn/third_party/xbyak_aarch64/src/util_impl_linux.h, 444: Can't open MIDR_EL1 sysfs entry1062xFormers 0.0.351063memory_efficient_attention.ckF: unavailable1064memory_efficient_attention.ckB: unavailable1065memory_efficient_attention.ck_splitKF: unavailable1066memory_efficient_attention.cutlassF-pt: available1067memory_efficient_attention.cutlassB-pt: available1068memory_efficient_attention.cutlassF-blackwell: unavailable1069memory_efficient_attention.cutlassB-blackwell: unavailable1070memory_efficient_attention.fa2F@2.5.7-pt: available1071memory_efficient_attention.fa2B@2.5.7-pt: available1072memory_efficient_attention.fa3F@0.0.0: unavailable1073memory_efficient_attention.fa3B@0.0.0: unavailable1074memory_efficient_attention.fa3F_splitKV@0.0.0: unavailable1075memory_efficient_attention.triton_splitKF: available1076indexing.scaled_index_addF: unavailable1077indexing.scaled_index_addB: unavailable1078indexing.index_select: unavailable1079sp24.sparse24_sparsify_both_ways: available1080sp24.sparse24_apply: available1081sp24.sparse24_apply_dense_output: available1082sp24._sparse24_gemm: available1083sp24._cslt_sparse_mm_search@0.8.1: available1084sp24._cslt_sparse_mm@0.8.1: available1085is_triton_available: False1086pytorch.version: 2.13.01087pytorch.cuda: not available1088dcgm_profiler: unavailable1089build.info: available1090build.cuda_version: 13021091build.hip_version: None1092build.python_version: 3.13.151093build.torch_version: 2.13.01094build.env.TORCH_CUDA_ARCH_LIST: 12.0;12.11095build.env.PYTORCH_ROCM_ARCH: None1096build.env.XFORMERS_BUILD_TYPE: None1097build.env.XFORMERS_ENABLE_DEBUG_ASSERTIONS: None1098build.env.NVCC_FLAGS: None1099build.env.XFORMERS_PACKAGE_FROM: None1100build.nvcc_version: 13.2.861101source.privacy: open source1102setting known thread variables to default: 11103 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=11104pytest flags: -m pytest tests --ignore-glob=tests/test_fwbw_overlap.py1105============================= test session starts ==============================1106platform linux -- Python 3.13.15, pytest-9.1.1, pluggy-1.6.01107rootdir: /build/source1108configfile: pyproject.toml1109plugins: hydra-core-1.3.5, cov-stub-1.0.0, timeout-2.4.0, anyio-4.14.21110collected 3829 items 11111112tests/test_attention_patterns.py ....................................... [ 1%]1113........................................................................ [ 2%]1114......................... [ 3%]1115tests/test_checkpoint.py ................................sssssssssssssss [ 4%]1116sssssssssssssssss........ssssssssssssssssssssssssssssssssssssssss [ 6%]1117tests/test_fmha_flop_formula.py ssssssssssssssss. [ 6%]1118tests/test_fmha_merge_attentions.py ssssssssssssssssssssssssssssssssssss [ 7%]1119ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%]1120sssssssssss [ 10%]1121tests/test_indexing.py sssssss [ 10%]1122tests/test_mem_eff_attention.py ssssssssssssssssssssssssssssssssssssssss [ 11%]1123ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 13%]1124ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 15%]1125ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 16%]1126ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 18%]1127ssssssssssssss.......sssssssssssssssssssssssssssssssssssssssssssssssssss [ 20%]1128ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 22%]1129ssssssssssssssssssssssssssssssssssssss.ssssss.sssssss..sssssssssssssssss [ 24%]1130ssssssssssssssssssssssssssssssssssssss.sssssssssssssssssssssssssssssssss [ 26%]1131ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 28%]1132ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 30%]1133ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 31%]1134ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 33%]1135ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 35%]1136ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 37%]1137ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 39%]1138ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 41%]1139ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 43%]1140ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 45%]1141ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 46%]1142ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 48%]1143ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 50%]1144ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 52%]1145ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 54%]1146ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 56%]1147ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 58%]1148ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 60%]1149ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 62%]1150ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 63%]1151ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 65%]1152ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 67%]1153ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 69%]1154ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 71%]1155ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 73%]1156ssssssssssssssss [ 73%]1157tests/test_multiprocessing_utils.py . [ 73%]1158tests/test_profiler.py ssssssssssssssss [ 74%]1159tests/test_rmsnorm.py ssssssssssssssssss [ 74%]1160tests/test_rope_padded.py ssssssssssssssssssssssssssssssssssssssssssssss [ 75%]1161ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 77%]1162ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 79%]1163ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 81%]1164sssssssssssssssssssssssssssss [ 82%]1165tests/test_seqpar.py sssssssssssssssssssssssssssssssssssssssssssssssssss [ 83%]1166sssssssssssssssssssssssssssssssssssssssssssss [ 84%]1167tests/test_sequence_parallel_fused_ops.py ssssssssssssssssssssssssssssss [ 85%]1168ssssssssssssssssssssssssssssssssssssssssssssss [ 86%]1169tests/test_sparse_tensors.py ..... [ 86%]1170tests/test_sparsity24.py sssssssssssssssssssssssssssssssssssssssssssssss [ 88%]1171sssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 89%]1172tests/test_splitk_reference.py .................. [ 90%]1173tests/test_tiled_matmul.py sssssssssssssssssssssssssssssssssssssssssssss [ 91%]1174sssssssssssssss [ 91%]1175tests/test_tree_attention.py sssssssssssssssssssssssssssssssssssssssssss [ 92%]1176ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 94%]1177ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 96%]1178ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 98%]1179sssssssssssss.................... [ 99%]1180tests/test_triton_varargs.py sssss [ 99%]1181tests/test_unbind.py .................... [100%]11821183=============================== warnings summary ===============================1184tests/test_attention_patterns.py::test_local_1d_pattern[50-3]1185 /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.)1186 return _VF.meshgrid(tensors, **kwargs) # type: ignore[attr-defined]11871188tests/test_checkpoint.py: 16 warnings1189 /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).1190 out = checkpoint(module[i], out, policy_fn=policy_fn)11911192tests/test_checkpoint.py: 16 warnings1193 /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).1194 out = checkpoint(module, out, policy_fn=policy_fn)11951196tests/test_checkpoint.py::test_list_operators1197 /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.1198 operators = list_operators(module, inputs)11991200tests/test_tree_attention.py: 16 warnings1201 /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.1202 torch.testing.assert_allclose(12031204-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html1205=============== 253 passed, 3576 skipped, 50 warnings in 15.04s ================1206Finished executing pytestCheckPhase