FBGEMM GenAI support for SM100 & SM120

[flaviusburca]

Currently the build fails on fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu:

FAILED: [code=255] experimental/gen_ai/CMakeFiles/fbgemm_gpu_experimental_gen_ai.dir/src/quantize/quantize.cu.o 
/usr/local/cuda-12.9/bin/nvcc -forward-unknown-to-host-compiler -DUSE_C10D_GLOO -DUSE_C10D_MPI -DUSE_C10D_NCCL -DUSE_DISTRIBUTED -DUSE_RPC -DUSE_TENSORPIPE -Dfbgemm_gpu_experimental_gen_ai_EXPORTS -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../include -I/opt/densemax/train/FBGEMM/fbgemm_gpu -I/opt/densemax/train/FBGEMM/fbgemm_gpu/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/asmjit/src -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/cpuinfo/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/cutlass/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/cutlass/tools/util/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/composable_kernel/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/composable_kernel/library/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/../external/json/include -I/usr/lib/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/experimental/gen_ai/src/quantize -I/opt/densemax/train/FBGEMM/fbgemm_gpu/experimental/gen_ai/src/quantize/common/include -I/opt/densemax/train/FBGEMM/fbgemm_gpu/experimental/gen_ai/src/kv_cache -I/opt/densemax/train/FBGEMM/fbgemm_gpu/experimental/gen_ai/../../include -isystem /opt/miniforge/envs/train/lib/python3.12/site-packages/torch/include -isystem /opt/miniforge/envs/train/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -I/usr/lib/include -isystem /opt/miniforge/envs/train/lib/python3.12/site-packages/torch/include -isystem /opt/miniforge/envs/train/lib/python3.12/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda-12.9/include -DONNX_NAMESPACE=onnx_c2 -gencode arch=compute_120,code=sm_120 -Xcudafe --diag_suppress=cc_clobber_ignored,--diag_suppress=field_without_dll_interface,--diag_suppress=base_class_has_different_dll_interface,--diag_suppress=dll_interface_conflict_none_assumed,--diag_suppress=dll_interface_conflict_dllexport_assumed,--diag_suppress=bad_friend_decl --expt-relaxed-constexpr --expt-extended-lambda -O3 -DNDEBUG -std=c++20 -Xcompiler=-fPIC -Wno-deprecated-enum-enum-conversion -Wno-deprecated-declarations -MD -MT experimental/gen_ai/CMakeFiles/fbgemm_gpu_experimental_gen_ai.dir/src/quantize/quantize.cu.o -MF experimental/gen_ai/CMakeFiles/fbgemm_gpu_experimental_gen_ai.dir/src/quantize/quantize.cu.o.d -x cu -c /opt/densemax/train/FBGEMM/fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu -o experimental/gen_ai/CMakeFiles/fbgemm_gpu_experimental_gen_ai.dir/src/quantize/quantize.cu.o
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 4969; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 4970; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 4971; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 4972; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5185; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5186; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5187; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5188; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5369; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5370; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5371; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5372; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5671; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5672; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5673; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5674; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5896; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5897; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5898; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 5899; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6089; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6090; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6091; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6092; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6370; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6371; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6372; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6373; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6574; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6575; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6576; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6577; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6746; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6747; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6748; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 6749; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7036; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7037; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7038; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7039; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7249; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7250; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7251; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7252; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7430; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7431; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7432; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas /tmp/tmpxft_00000c40_00000000-6_quantize.ptx, line 7433; error   : Instruction 'cvt with .e2m1x2' not supported on .target 'sm_120'
ptxas fatal   : Ptx assembly aborted due to errors

[q10]

Hi @flaviusburca, unfortunately we only support sm100a (i.e. no sm100).

Please see the FBGEMM releases compatibility table here.