[FEA] Support of FP32 scales in FP8 scaledMM

[sanchitintel]

Summary

Some models use Float scales even for FP8 weights. FP8 ScaledMM should be able to support FP32 weights

Details

Currently, FP16 scales are being used in FP8 scaledMM.
I tried to use FP32 scales by using an FP32 copy-atom for scales, which is similar to the FP16 copy-atom being used to copy scales, but it didn't help

https://github.com/codeplaysoftware/cutlass-sycl/compare/sycl-develop...sanchitintel:enable_float_scales_fp8_scaledmm?expand=1

Please advise how to go about adding support for Float32 scales in FP8 scaledMM.

Thank you!

[sanchitintel]

The build with https://github.com/codeplaysoftware/cutlass-sycl/compare/sycl-develop...sanchitintel:enable_float_scales_fp8_scaledmm?expand=1 succeeded at my end when I rebased to the latest sycl-develop branch.

Sorry for the false alarm!

[sanchitintel]

Sorry @aacostadiaz, the branch I had created didn't have all the changes. I pushed all changes now. This issue can be reproduced using the branch I created for debugging.

If possible, can you please take a look? Thanks!

Basically, by changing scales and Zero dtype to float, compilation is failing after adding a scale_zero_copy_traits for FP32 that's similar to that for FP16 or int8.

[1/5] Building CXX object examples/sycl/08_bmg_gemm_f8/CMakeFiles/08_bmg_gemm_f8_scaling.dir/08_bmg_gemm_f8_scaling.cpp.o
FAILED: examples/sycl/08_bmg_gemm_f8/CMakeFiles/08_bmg_gemm_f8_scaling.dir/08_bmg_gemm_f8_scaling.cpp.o
/localdiskj/dpcpp/bin/clang++ -DCUTLASS_ENABLE_SYCL -DCUTLASS_SYCL_BUILTIN_ENABLE -DCUTLASS_SYCL_PROFILING_ENABLED -DSYCLCOMPAT_PROFILING_ENABLED -DSYCL_INTEL_TARGET -I/localdiskj/scaled/include -I/localdiskj/scaled/examples/common -I/localdiskj/scaled/build/deps/oneMKL/include -I/localdiskj/scaled/build/include -I/localdiskj/scaled/tools/util/include -isystem /include -isystem /include/cccl -ftemplate-backtrace-limit=0 -fdiagnostics-color=always -DCUTLASS_VERSIONS_GENERATED -Wno-error=implicit-int-conversion -Wno-error=pass-failed -Wno-error=inconsistent-missing-override -Wno-sign-conversion -Wno-unused-parameter -O2 -g -DNDEBUG -std=c++17 -fPIE -fsycl -fsycl-targets=intel_gpu_pvc -Xspirv-translator -spirv-ext=+SPV_INTEL_split_barrier -fno-sycl-instrument-device-code -MD -MT examples/sycl/08_bmg_gemm_f8/CMakeFiles/08_bmg_gemm_f8_scaling.dir/08_bmg_gemm_f8_scaling.cpp.o -MF examples/sycl/08_bmg_gemm_f8/CMakeFiles/08_bmg_gemm_f8_scaling.dir/08_bmg_gemm_f8_scaling.cpp.o.d -o examples/sycl/08_bmg_gemm_f8/CMakeFiles/08_bmg_gemm_f8_scaling.dir/08_bmg_gemm_f8_scaling.cpp.o -c /localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp
clang++: warning: argument unused during compilation: '-Xspirv-translator -spirv-ext=+SPV_INTEL_split_barrier' [-Wunused-command-line-argument]
In file included from /localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:50:
In file included from /localdiskj/scaled/include/cutlass/epilogue/collective/default_epilogue.hpp:39:
In file included from /localdiskj/scaled/include/cutlass/gemm/dispatch_policy.hpp:38:
In file included from /localdiskj/scaled/include/cute/atom/copy_traits_sm100.hpp:38:
In file included from /localdiskj/scaled/include/cute/atom/copy_atom.hpp:761:
In file included from /localdiskj/scaled/include/cute/atom/copy_traits_sm100_tma.hpp:40:
In file included from /localdiskj/scaled/include/cute/tensor.hpp:51:
/localdiskj/scaled/include/cute/algorithm/copy.hpp:178:17: error: static assertion failed due to requirement 'Layout<cute::C<2>, cute::ScaledBasis<cute::C<16>, 0>>::rank == Layout<cute::tuple<cute::C<1>, cute::C<2>>, cute::tuple<cute::C<0>, cute::C<1>>>::rank': CopyAtom rank-mismatch.
  178 |   static_assert(SrcLayout::rank == DstLayout::rank, "CopyAtom rank-mismatch.");
      |                 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/localdiskj/scaled/include/cute/algorithm/copy.hpp:514:10: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::C<2>, cute::ScaledBasis<cute::C<16>, 0>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::C<1>, cute::C<2>>>>' requested here
  514 |   return copy(copy_policy, src, dst);
      |          ^
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:109:14: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::C<2>, cute::ScaledBasis<cute::C<16>, 0>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::C<1>, cute::C<2>>>>' requested here
  109 |       return copy(*this, tensor<0>(src), tensor<0>(dst));
      |              ^
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:124:12: note: in instantiation of function template specialization 'cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>::call<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>>>>' requested here
  124 |     return call(src, dst);
      |            ^
/localdiskj/scaled/include/cute/algorithm/copy.hpp:211:19: note: in instantiation of function template specialization 'cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>::call<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>>>>' requested here
  211 |         copy_atom.call(src_c(_,i), dst_c(_,i));
      |                   ^
/localdiskj/scaled/include/cute/algorithm/copy.hpp:412:10: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>, cute::C<1>, cute::C<1>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>, cute::ScaledBasis<cute::C<32>, 0>, cute::C<0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>, cute::C<1>, cute::C<1>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::C<0>, cute::C<0>>>>' requested here
  412 |   return copy(static_cast<CopyAtom const&>(tiled_copy), src, dst);
      |          ^
/localdiskj/scaled/include/cutlass/gemm/collective/xe_mma_fp8_scaling.hpp:648:9: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>, cute::Layout<cute::tuple<cute::C<16>, cute::C<1>>>, cute::tuple<cute::C<16>, cute::C<1>>, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>, cute::C<1>, cute::C<1>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>, cute::ScaledBasis<cute::C<32>, 0>, cute::C<0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>, cute::C<1>, cute::C<1>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::C<0>, cute::C<0>>>>' requested here
  648 |         copy(mainloop.tiled_copy_scaleA, copy_iter_sA(_, _, _, k_tile / k_reload_factor), copy_tCrSA);
      |         ^
/localdiskj/scaled/include/cutlass/gemm/kernel/xe_gemm.hpp:258:5: note: in instantiation of function template specialization 'cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>::operator()<cute::Tensor<cute::ArrayEngine<float, 128>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<8>, cute::C<32>>>>, cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, cute::C<0>, int>>>, cute::Layout<cute::tuple<cute::C<256>, cute::C<32>, int>, cute::tuple<cute::ScaledBasis<cute::C<1>, 0>, cute::ScaledBasis<cute::C<1>, 1>, cute::ScaledBasis<cute::C<32>, 1>>>>, cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, cute::C<0>, int>>>, cute::Layout<cute::tuple<cute::C<256>, cute::C<32>, int>, cute::tuple<cute::ScaledBasis<cute::C<1>, 0>, cute::ScaledBasis<cute::C<1>, 1>, cute::ScaledBasis<cute::C<32>, 1>>>>, cute::Tensor<cute::ArrayEngine<float, 128>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<8>, cute::C<32>>>>, cute::ForwardCoordIterator<cute::tuple<int>, cute::tuple<int>, cute::tuple<cute::ScaledBasis<cute::C<1>, 0>>>, cute::tuple<int, int, cute::Underscore, int>>' requested here
  258 |     collective_mma(
      |     ^
/localdiskj/scaled/include/cutlass/device_kernel.h:180:3: note: in instantiation of member function 'cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>::operator()' requested here
  180 |   op(params, smem);
      |   ^
/localdiskj/scaled/include/cutlass/gemm/device/gemm_universal_adapter.h:589:55: note: in instantiation of function template specialization 'cutlass::device_kernel<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>' requested here
  589 |         auto event = syclcompat::experimental::launch<device_kernel<GemmKernel>>(policy, q, params);
      |                                                       ^
/localdiskj/scaled/include/cutlass/gemm/device/gemm_universal_adapter.h:661:12: note: in instantiation of member function 'cutlass::gemm::device::GemmUniversalAdapter<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>::run' requested here
  661 |     return run(params_, stream, cuda_adapter, launch_with_pdl);
      |            ^
/localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:429:27: note: in instantiation of member function 'cutlass::gemm::device::GemmUniversalAdapter<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>::run' requested here
  429 |     CUTLASS_CHECK(gemm_op.run());
      |                           ^
/localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:567:41: note: in instantiation of member function 'ExampleRunner<cutlass::gemm::device::GemmUniversalAdapter<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>>::run' requested here
  567 |     CUTLASS_CHECK(ExampleRunner<Gemm>{}.run(options, hw_info));
      |                                         ^
/localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:590:3: note: in instantiation of function template specialization 'launcher<cutlass::float_e5m2_t>' requested here
  590 |   launcher<cutlass::float_e5m2_t>(options);
      |   ^
/localdiskj/scaled/include/cute/algorithm/copy.hpp:178:33: note: expression evaluates to '1 == 2'
  178 |   static_assert(SrcLayout::rank == DstLayout::rank, "CopyAtom rank-mismatch.");
      |                 ~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~
In file included from /localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:50:
In file included from /localdiskj/scaled/include/cutlass/epilogue/collective/default_epilogue.hpp:39:
In file included from /localdiskj/scaled/include/cutlass/gemm/dispatch_policy.hpp:38:
In file included from /localdiskj/scaled/include/cute/atom/copy_traits_sm100.hpp:38:
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:98:19: error: static assertion failed due to requirement 'Layout<cute::tuple<cute::C<1>, cute::C<2>>, cute::tuple<cute::C<0>, cute::C<1>>>::rank == 1': Expected rank-1 dst tensor
   98 |     static_assert(DLayout::rank == 1, "Expected rank-1 dst tensor");
      |                   ^~~~~~~~~~~~~~~~~~
/localdiskj/scaled/include/cute/algorithm/copy.hpp:181:15: note: in instantiation of function template specialization 'cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>::call<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::C<2>, cute::ScaledBasis<cute::C<16>, 0>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::C<1>, cute::C<2>>>>' requested here
  181 |     copy_atom.call(src, dst);
      |               ^
/localdiskj/scaled/include/cute/algorithm/copy.hpp:514:10: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::C<2>, cute::ScaledBasis<cute::C<16>, 0>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::C<1>, cute::C<2>>>>' requested here
  514 |   return copy(copy_policy, src, dst);
      |          ^
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:109:14: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::C<2>, cute::ScaledBasis<cute::C<16>, 0>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::C<1>, cute::C<2>>>>' requested here
  109 |       return copy(*this, tensor<0>(src), tensor<0>(dst));
      |              ^
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:124:12: note: in instantiation of function template specialization 'cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>::call<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>>>>' requested here
  124 |     return call(src, dst);
      |            ^
/localdiskj/scaled/include/cute/algorithm/copy.hpp:211:19: note: in instantiation of function template specialization 'cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>::call<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>>>>' requested here
  211 |         copy_atom.call(src_c(_,i), dst_c(_,i));
      |                   ^
/localdiskj/scaled/include/cute/algorithm/copy.hpp:412:10: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>, cute::C<1>, cute::C<1>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>, cute::ScaledBasis<cute::C<32>, 0>, cute::C<0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>, cute::C<1>, cute::C<1>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::C<0>, cute::C<0>>>>' requested here
  412 |   return copy(static_cast<CopyAtom const&>(tiled_copy), src, dst);
      |          ^
/localdiskj/scaled/include/cutlass/gemm/collective/xe_mma_fp8_scaling.hpp:648:9: note: in instantiation of function template specialization 'cute::copy<cute::Copy_Atom<cute::Copy_Traits<cute::XE_2D_U32x1x16_LD_N, cute::tuple<cute::C<1>, long, long>>, float>, cute::Layout<cute::tuple<cute::C<16>, cute::C<1>>>, cute::tuple<cute::C<16>, cute::C<1>>, cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>, cute::Layout<cute::tuple<cute::C<2>, cute::C<1>, cute::C<1>>, cute::tuple<cute::ScaledBasis<cute::C<16>, 0>, cute::ScaledBasis<cute::C<32>, 0>, cute::C<0>>>, cute::ViewEngine<float *>, cute::Layout<cute::tuple<cute::tuple<cute::C<1>, cute::C<2>>, cute::C<1>, cute::C<1>>, cute::tuple<cute::tuple<cute::C<0>, cute::C<1>>, cute::C<0>, cute::C<0>>>>' requested here
  648 |         copy(mainloop.tiled_copy_scaleA, copy_iter_sA(_, _, _, k_tile / k_reload_factor), copy_tCrSA);
      |         ^
/localdiskj/scaled/include/cutlass/gemm/kernel/xe_gemm.hpp:258:5: note: in instantiation of function template specialization 'cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>::operator()<cute::Tensor<cute::ArrayEngine<float, 128>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<8>, cute::C<32>>>>, cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, cute::C<0>, int>>>, cute::Layout<cute::tuple<cute::C<256>, cute::C<32>, int>, cute::tuple<cute::ScaledBasis<cute::C<1>, 0>, cute::ScaledBasis<cute::C<1>, 1>, cute::ScaledBasis<cute::C<32>, 1>>>>, cute::Tensor<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, cute::C<0>, int>>>, cute::Layout<cute::tuple<cute::C<256>, cute::C<32>, int>, cute::tuple<cute::ScaledBasis<cute::C<1>, 0>, cute::ScaledBasis<cute::C<1>, 1>, cute::ScaledBasis<cute::C<32>, 1>>>>, cute::Tensor<cute::ArrayEngine<float, 128>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<8>, cute::C<32>>>>, cute::ForwardCoordIterator<cute::tuple<int>, cute::tuple<int>, cute::tuple<cute::ScaledBasis<cute::C<1>, 0>>>, cute::tuple<int, int, cute::Underscore, int>>' requested here
  258 |     collective_mma(
      |     ^
/localdiskj/scaled/include/cutlass/device_kernel.h:180:3: note: in instantiation of member function 'cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>::operator()' requested here
  180 |   op(params, smem);
      |   ^
/localdiskj/scaled/include/cutlass/gemm/device/gemm_universal_adapter.h:589:55: note: in instantiation of function template specialization 'cutlass::device_kernel<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>' requested here
  589 |         auto event = syclcompat::experimental::launch<device_kernel<GemmKernel>>(policy, q, params);
      |                                                       ^
/localdiskj/scaled/include/cutlass/gemm/device/gemm_universal_adapter.h:661:12: note: in instantiation of member function 'cutlass::gemm::device::GemmUniversalAdapter<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>::run' requested here
  661 |     return run(params_, stream, cuda_adapter, launch_with_pdl);
      |            ^
/localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:429:27: note: in instantiation of member function 'cutlass::gemm::device::GemmUniversalAdapter<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>::run' requested here
  429 |     CUTLASS_CHECK(gemm_op.run());
      |                           ^
/localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:567:41: note: in instantiation of member function 'ExampleRunner<cutlass::gemm::device::GemmUniversalAdapter<cutlass::gemm::kernel::GemmUniversal<cute::tuple<int, int, int, int>, cutlass::gemm::collective::CollectiveMma<cutlass::gemm::MainloopIntelXeXMX16FP8Scaling<2>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<long, cute::C<1>, long>, cute::tuple<cutlass::float_e5m2_t, float, cute::tuple<cute::C<1>, long, long>>, cute::tuple<cute::C<1>, long, long>, cute::TiledMMA<cute::MMA_Atom<cute::XE_8x16x16_F32F16F16F32_TT>, cute::Layout<cute::tuple<cute::C<8>, cute::C<4>, cute::C<1>>, cute::tuple<cute::C<4>, cute::C<1>, cute::C<0>>>, const cute::tuple<cute::Layout<cute::tuple<cute::C<8>, cute::C<8>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<32>, cute::C<8>>>, cute::Layout<cute::tuple<cute::C<16>, cute::C<4>, cute::C<4>>, cute::tuple<cute::C<1>, cute::C<64>, cute::C<16>>>, cute::Layout<cute::C<32>>>>, cute::XE_2D_U8x32x32_LD_N, void, void, cute::identity, cute::XE_2D_U8x32x32_LD_V, void, void, cute::identity>, cutlass::epilogue::collective::CollectiveEpilogue<cutlass::epilogue::IntelXeXMX16, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, float, cute::tuple<long, cute::C<1>, long>, float, cute::tuple<long, cute::C<1>, long>, cutlass::epilogue::fusion::FusionCallbacks<cutlass::epilogue::IntelXeXMX16, cutlass::epilogue::fusion::LinearCombination<float, float, float, float, cutlass::FloatRoundStyle::round_to_nearest>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>, cute::tuple<cute::C<256>, cute::C<256>, cute::C<32>>>, cute::XE_2D_U32x8x16_LD_N, void, void, cute::XE_2D_U32x8x16_ST_N, void, void>>>>::run' requested here
  567 |     CUTLASS_CHECK(ExampleRunner<Gemm>{}.run(options, hw_info));
      |                                         ^
/localdiskj/scaled/examples/sycl/08_bmg_gemm_f8/08_bmg_gemm_f8_scaling.cpp:590:3: note: in instantiation of function template specialization 'launcher<cutlass::float_e5m2_t>' requested here
  590 |   launcher<cutlass::float_e5m2_t>(options);
      |   ^
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:98:33: note: expression evaluates to '2 == 1'
   98 |     static_assert(DLayout::rank == 1, "Expected rank-1 dst tensor");
      |                   ~~~~~~~~~~~~~~^~~~
/localdiskj/scaled/include/cute/atom/copy_atom.hpp:111:21: error: static assertion failed due to requirement 'dependent_false<cute::ViewEngine<cute::ArithmeticTupleIterator<cute::ArithmeticTuple<int, int, int>>>>': CopyAtom: Src/Dst partitioning does not match the instruction requirement.
  111 |       static_assert(dependent_false<SEngine>,
      |                     ^~~~~~~~~~~~~~~~~~~~~~~~
3 errors generated.