NVIDIA / NVIDIA/cutlass

[BUG] Some CUTLASS headers are not self-contained

Open
#1,827 9 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage bug
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Describe the bug
Some files are missing the headers that they rely on, which means they cannot be included by themselves. This is "hidden" in most of the examples because they import many things and one of the other headers will often include what's needed, but this is fragile and has caused build failures as we rearrange includes. I assume everything in cutlass/include and cutlass/tools/util/include/ are supposed to be includable individually? If so, then the following headers are broken (generated with a simple shell script):

Broken headers
cute/algorithm/cooperative_copy.hpp
cute/algorithm/copy.hpp
cute/algorithm/functional.hpp
cute/algorithm/gemm.hpp
cute/algorithm/prefetch.hpp
cute/arch/mma_sm90.hpp
cute/atom/copy_atom.hpp
cute/atom/copy_traits_sm90_tma.hpp
cute/atom/mma_traits_sm90.hpp
cute/atom/mma_traits_sm90_gmma.hpp
cute/layout_composed.hpp
cute/numeric/int.hpp
cute/numeric/numeric_types.hpp
cute/pointer_flagged.hpp
cutlass/arch/mma_sm61.h
cutlass/arch/mma_sm89.h
cutlass/arch/reg_reconfig.h
cutlass/arch/simd.h
cutlass/arch/simd_sm60.h
cutlass/arch/simd_sm61.h
cutlass/arch/wmma_sm70.h
cutlass/arch/wmma_sm72.h
cutlass/arch/wmma_sm75.h
cutlass/conv/collective/builders/sm90_gmma_builder.inl
cutlass/conv/collective/collective_builder.hpp
cutlass/conv/collective/collective_conv.hpp
cutlass/conv/collective/detail.hpp
cutlass/conv/collective/sm90_implicit_gemm_gmma_ss_warpspecialized.hpp
cutlass/conv/convnd_problem_shape.hpp
cutlass/conv/device/implicit_gemm_convolution.h
cutlass/conv/device/implicit_gemm_convolution_fusion.h
cutlass/conv/kernel/implicit_gemm_convolution_strided_dgrad.h
cutlass/conv/kernel/sm90_implicit_gemm_tma_warpspecialized.hpp
cutlass/conv/threadblock/conv2d_tile_iterator.h
cutlass/conv/threadblock/conv2d_wgrad_activation_tile_access_iterator_optimized.h
cutlass/conv/threadblock/conv2d_wgrad_output_gradient_tile_access_iterator_optimized.h
cutlass/conv/threadblock/depthwise_fprop_direct_conv_multistage.h
cutlass/conv/threadblock/depthwise_fprop_filter_tile_access_iterator_direct_conv_optimized.h
cutlass/conv/threadblock/depthwise_mma_base.h
cutlass/conv/threadblock/depthwise_mma_core_with_lane_access_size.h
cutlass/conv/threadblock/implicit_gemm_multistage.h
cutlass/conv/threadblock/predicated_scale_bias_vector_access_iterator.h
cutlass/conv/threadblock/predicated_scale_bias_vector_iterator.h
cutlass/device_kernel.h
cutlass/epilogue/collective/builders/sm90_builder.inl
cutlass/epilogue/collective/builders/sm90_common.inl
cutlass/epilogue/collective/collective_builder.hpp
cutlass/epilogue/collective/collective_epilogue.hpp
cutlass/epilogue/collective/detail.hpp
cutlass/epilogue/collective/epilogue_tensor_broadcast.hpp
cutlass/epilogue/collective/sm90_epilogue_array_tma_warpspecialized.hpp
cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized.hpp
cutlass/epilogue/collective/sm90_epilogue_tma_warpspecialized_bias_elementwise.hpp
cutlass/epilogue/fusion/callbacks.hpp
cutlass/epilogue/fusion/operations.hpp
cutlass/epilogue/fusion/sm90_callbacks_tma_warpspecialized.hpp
cutlass/epilogue/fusion/sm90_visitor_load_tma_warpspecialized.hpp
cutlass/epilogue/fusion/sm90_visitor_store_tma_warpspecialized.hpp
cutlass/epilogue/fusion/sm90_visitor_tma_warpspecialized.hpp
cutlass/epilogue/thread/linear_combination_params.h
cutlass/epilogue/thread/linear_combination_tensor_broadcast.hpp
cutlass/epilogue/threadblock/default_epilogue_complex_tensor_op.h
cutlass/epilogue/threadblock/default_epilogue_complex_tensor_op_blas3.h
cutlass/epilogue/threadblock/default_epilogue_direct_store.h
cutlass/epilogue/threadblock/default_epilogue_tensor_op_blas3.h
cutlass/epilogue/threadblock/default_epilogue_with_broadcast.h
cutlass/epilogue/threadblock/default_epilogue_wmma_tensor_op.h
cutlass/epilogue/threadblock/epilogue_base_streamk.h
cutlass/epilogue/threadblock/epilogue_depthwise.h
cutlass/epilogue/threadblock/epilogue_direct_store.h
cutlass/epilogue/threadblock/epilogue_gemm_k_reduction.h
cutlass/epilogue/threadblock/epilogue_smem_accumulator.h
cutlass/epilogue/threadblock/epilogue_streamk_with_broadcast.h
cutlass/epilogue/threadblock/epilogue_visitor_with_softmax.h
cutlass/epilogue/threadblock/epilogue_with_visitor_callbacks.h
cutlass/epilogue/threadblock/epilogue_workspace.h
cutlass/epilogue/threadblock/fusion/visitor_2x.hpp
cutlass/epilogue/threadblock/fusion/visitor_compute.hpp
cutlass/epilogue/threadblock/fusion/visitor_load.hpp
cutlass/epilogue/threadblock/fusion/visitor_store.hpp
cutlass/epilogue/threadblock/fusion/visitors.hpp
cutlass/epilogue/threadblock/output_tile_thread_map.h
cutlass/epilogue/threadblock/predicated_tile_iterator_blas3.h
cutlass/epilogue/threadblock/shared_load_iterator.h
cutlass/epilogue/threadblock/shared_load_iterator_mixed.h
cutlass/epilogue/threadblock/shared_load_iterator_pitch_linear.h
cutlass/epilogue/warp/fragment_iterator_complex_tensor_op.h
cutlass/epilogue/warp/fragment_iterator_gaussian_complex_tensor_op.h
cutlass/epilogue/warp/fragment_iterator_wmma_tensor_op.h
cutlass/epilogue/warp/tile_iterator_simt.h
cutlass/epilogue/warp/tile_iterator_tensor_op_mixed.h
cutlass/epilogue/warp/tile_iterator_volta_tensor_op.h
cutlass/epilogue/warp/tile_iterator_wmma_tensor_op.h
cutlass/floating_point_nvrtc.h
cutlass/gemm/collective/fp8_accumulation.hpp
cutlass/gemm/collective/sm80_mma_multistage.hpp
cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized.hpp
cutlass/gemm/collective/sm90_mma_multistage_gmma_rs_warpspecialized.hpp
cutlass/gemm/collective/sm90_mma_multistage_gmma_ss_warpspecialized.hpp
cutlass/gemm/collective/sm90_mma_tma_gmma_rs_warpspecialized.hpp
cutlass/gemm/collective/sm90_mma_tma_gmma_rs_warpspecialized_mixed_input.hpp
cutlass/gemm/collective/sm90_mma_tma_gmma_ss.hpp
cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized.hpp
cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8.hpp
cutlass/gemm/device/rank_2k_grouped.h
cutlass/gemm/kernel/default_gemm_streamk_with_broadcast.h
cutlass/gemm/kernel/default_rank_2k_grouped.h
cutlass/gemm/kernel/gemm_params.h
cutlass/gemm/kernel/gemm_streamk_with_fused_epilogue.h
cutlass/gemm/kernel/params_sparse_base.h
cutlass/gemm/kernel/sm70_gemm.hpp
cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized.hpp
cutlass/gemm/kernel/sm90_gemm_warpspecialized.hpp
cutlass/gemm/kernel/sm90_gemm_warpspecialized_cooperative.hpp
cutlass/gemm/kernel/sm90_tile_scheduler_group.hpp
cutlass/gemm/thread/mma_sm50.h
cutlass/gemm/threadblock/default_ell_mma.h
cutlass/gemm/threadblock/default_mma_core.h
cutlass/gemm/threadblock/default_mma_core_sm70.h
cutlass/gemm/threadblock/default_mma_core_with_access_size.h
cutlass/gemm/threadblock/default_mma_core_wmma.h
cutlass/gemm/threadblock/ell_mma_multistage.h
cutlass/gemm/threadblock/ell_mma_pipelined.h
cutlass/gemm/threadblock/mma_blas3_multistage.h
cutlass/gemm/threadblock/mma_multistage.h
cutlass/gemm/threadblock/mma_planar_complex_base.h
cutlass/gemm/threadblock/mma_planar_complex_multistage.h
cutlass/gemm/threadblock/mma_planar_complex_pipelined.h
cutlass/gemm/threadblock/mma_singlestage.h
cutlass/gemm/threadblock/mma_sparse_base.h
cutlass/gemm/threadblock/mma_sparse_multistage.h
cutlass/gemm/threadblock/mma_with_reduction_multistage.h
cutlass/layout/tensor_op_multiplicand_sm70.h
cutlass/pipeline/pipeline.hpp
cutlass/pipeline/sm90_pipeline.hpp
cutlass/real.h
cutlass/reduction/device/reduce_split_k.h
cutlass/reduction/kernel/reduce_softmax_final.h
cutlass/reduction/threadblock_swizzle.h
cutlass/tensor_view_planar_complex.h
cutlass/tfloat32.h
cutlass/thread/matrix.h
cutlass/transform/collective/sm90_wgmma_transpose.hpp
cutlass/transform/thread/transpose.h
cutlass/transform/threadblock/ell_iterator.h
cutlass/transform/threadblock/predicated_scale_bias_vector_iterator.h
cutlass/transform/threadblock/predicated_vector_access_iterator.h
cutlass/transform/threadblock/regular_tile_access_iterator.h
cutlass/transform/threadblock/regular_tile_iterator_pitch_linear_2dthreadtile.h
cutlass/transform/threadblock/vector_iterator.h
cutlass/util/cublas_wrappers.hpp
cutlass/util/device_rmsnorm.h
cutlass/util/host_tensor_planar_complex.h
cutlass/util/packed_stride.hpp
cutlass/util/reference/device/gemm.h
cutlass/util/reference/device/gett.hpp
cutlass/util/reference/device/kernel/gemm.h
cutlass/util/reference/device/kernel/tensor_elementwise.h
cutlass/util/reference/device/rank_2k_complex.h
cutlass/util/reference/device/thread/gemm.h
cutlass/util/reference/host/conv.hpp
cutlass/util/reference/host/error_metrics.h
cutlass/util/reference/host/symm_complex.h
cutlass/util/reference/host/tensor_copy.h
cutlass/util/reference/host/tensor_elementwise.h
cutlass/util/reference/host/tensor_foreach.h
cutlass/util/reference/host/tensor_norm.h
cutlass/util/reference/host/tensor_reduce.h

Steps/Code to reproduce bug
Here is an example:

#include <cutlass/epilogue/collective/collective_builder.hpp>

Build this with -c -std=c++20 -I cutlass/include/ -I cutlass/tools/util/include.

Expected behavior
This should compile cleanly.

Environment details (please complete the following information):

  • Environment location: Bare-metal
  • CUTLASS main branch
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Tue_Aug_15_22:02:13_PDT_2023
Cuda compilation tools, release 12.2, V12.2.140
Build cuda_12.2.r12.2/compiler.33191640_0

Additional context

Example output
$ nvcc -c -std=c++20 test.cu -I cutlass/include/ -I cutlass/tools/util/include/
cutlass/include/cutlass/epilogue/fusion/operations.hpp(134): error: identifier "sizeof_bits_v" is undefined
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                               ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(152): error: identifier "sizeof_bits_v" is undefined
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                               ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(169): error: identifier "sizeof_bits_v" is undefined
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                               ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(190): error: identifier "sizeof_bits_v" is undefined
    int AlignmentAux_ = 128 / sizeof_bits_v<ElementAux_>,
                              ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(191): error: type name is not allowed
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                             ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(191): error: expected a "," or ">"
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                                         ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(211): error: identifier "sizeof_bits_v" is undefined
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                               ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(212): error: type name is not allowed
    int AlignmentScalar_ = 128 / sizeof_bits_v<ElementScalar_>,
                                               ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(212): error: expected a "," or ">"
    int AlignmentScalar_ = 128 / sizeof_bits_v<ElementScalar_>,
                                                             ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(234): error: identifier "sizeof_bits_v" is undefined
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                               ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(264): error: identifier "sizeof_bits_v" is undefined
    int AlignmentAux_ = 128 / sizeof_bits_v<ElementAux_>,
                              ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(265): error: type name is not allowed
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                             ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(265): error: expected a "," or ">"
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                                         ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(291): error: identifier "sizeof_bits_v" is undefined
    int AlignmentAux_ = 128 / sizeof_bits_v<ElementAux_>,
                              ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(318): error: identifier "sizeof_bits_v" is undefined
    int AlignmentAux_ = 128 / sizeof_bits_v<ElementAux_>,
                              ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(319): error: type name is not allowed
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                             ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(319): error: expected a "," or ">"
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                                         ^

cutlass/include/cutlass/epilogue/collective/collective_builder.hpp(103): error: identifier "cute" is undefined
    cute::enable_if_t<not is_base_of_v<fusion::FusionOperation, FusionCallbacks>>
    ^

cutlass/include/cutlass/epilogue/collective/collective_builder.hpp(103): error: enable_if_t is not a template
    cute::enable_if_t<not is_base_of_v<fusion::FusionOperation, FusionCallbacks>>
          ^

cutlass/include/cutlass/epilogue/collective/collective_builder.hpp(103): error: identifier "is_base_of_v" is undefined
    cute::enable_if_t<not is_base_of_v<fusion::FusionOperation, FusionCallbacks>>
                          ^

cutlass/include/cutlass/epilogue/collective/collective_builder.hpp(103): error: expected a ">"
    cute::enable_if_t<not is_base_of_v<fusion::FusionOperation, FusionCallbacks>>
        ^

cutlass/include/cutlass/epilogue/collective/collective_builder.hpp(103): error: type name is not allowed
    cute::enable_if_t<not is_base_of_v<fusion::FusionOperation, FusionCallbacks>>
                                       ^

cutlass/include/cutlass/epilogue/collective/collective_builder.hpp(104): error: expected a ";"
  > {
  ^

cutlass/include/cute/arch/mma_sm90.hpp(154): error: complex is not a template
    using DRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(155): error: complex is not a template
    using ARegisters = complex<double>[2];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(156): error: complex is not a template
    using BRegisters = complex<double>[1];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(157): error: complex is not a template
    using CRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(160): error: complex is not a template
    fma(complex<double> & d0, complex<double> & d1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(160): error: complex is not a template
    fma(complex<double> & d0, complex<double> & d1,
                              ^

cutlass/include/cute/arch/mma_sm90.hpp(161): error: complex is not a template
        complex<double> & d2, complex<double> & d3,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(161): error: complex is not a template
        complex<double> & d2, complex<double> & d3,
                              ^

cutlass/include/cute/arch/mma_sm90.hpp(162): error: complex is not a template
        complex<double> const& a0, complex<double> const& a1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(162): error: complex is not a template
        complex<double> const& a0, complex<double> const& a1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(163): error: complex is not a template
        complex<double> const& b0,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(164): error: complex is not a template
        complex<double> const& c0, complex<double> const& c1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(164): error: complex is not a template
        complex<double> const& c0, complex<double> const& c1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(165): error: complex is not a template
        complex<double> const& c2, complex<double> const& c3)
        ^

cutlass/include/cute/arch/mma_sm90.hpp(165): error: complex is not a template
        complex<double> const& c2, complex<double> const& c3)
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(212): error: complex is not a template
    using DRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(213): error: complex is not a template
    using ARegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(214): error: complex is not a template
    using BRegisters = complex<double>[2];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(215): error: complex is not a template
    using CRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(218): error: complex is not a template
    fma(complex<double> & d0, complex<double> & d1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(218): error: complex is not a template
    fma(complex<double> & d0, complex<double> & d1,
                              ^

cutlass/include/cute/arch/mma_sm90.hpp(219): error: complex is not a template
        complex<double> & d2, complex<double> & d3,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(219): error: complex is not a template
        complex<double> & d2, complex<double> & d3,
                              ^

cutlass/include/cute/arch/mma_sm90.hpp(220): error: complex is not a template
        complex<double> const& a0, complex<double> const& a1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(220): error: complex is not a template
        complex<double> const& a0, complex<double> const& a1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(221): error: complex is not a template
        complex<double> const& a2, complex<double> const& a3,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(221): error: complex is not a template
        complex<double> const& a2, complex<double> const& a3,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(222): error: complex is not a template
        complex<double> const& b0, complex<double> const& b1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(222): error: complex is not a template
        complex<double> const& b0, complex<double> const& b1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(223): error: complex is not a template
        complex<double> const& c0, complex<double> const& c1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(223): error: complex is not a template
        complex<double> const& c0, complex<double> const& c1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(224): error: complex is not a template
        complex<double> const& c2, complex<double> const& c3)
        ^

cutlass/include/cute/arch/mma_sm90.hpp(224): error: complex is not a template
        complex<double> const& c2, complex<double> const& c3)
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(271): error: complex is not a template
    using DRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(272): error: complex is not a template
    using ARegisters = complex<double>[8];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(273): error: complex is not a template
    using BRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(274): error: complex is not a template
    using CRegisters = complex<double>[4];
                       ^

cutlass/include/cute/arch/mma_sm90.hpp(277): error: complex is not a template
    fma(complex<double> & d0, complex<double> & d1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(277): error: complex is not a template
    fma(complex<double> & d0, complex<double> & d1,
                              ^

cutlass/include/cute/arch/mma_sm90.hpp(278): error: complex is not a template
        complex<double> & d2, complex<double> & d3,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(278): error: complex is not a template
        complex<double> & d2, complex<double> & d3,
                              ^

cutlass/include/cute/arch/mma_sm90.hpp(279): error: complex is not a template
        complex<double> const& a0, complex<double> const& a1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(279): error: complex is not a template
        complex<double> const& a0, complex<double> const& a1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(280): error: complex is not a template
        complex<double> const& a2, complex<double> const& a3,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(280): error: complex is not a template
        complex<double> const& a2, complex<double> const& a3,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(281): error: complex is not a template
        complex<double> const& a4, complex<double> const& a5,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(281): error: complex is not a template
        complex<double> const& a4, complex<double> const& a5,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(282): error: complex is not a template
        complex<double> const& a6, complex<double> const& a7,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(282): error: complex is not a template
        complex<double> const& a6, complex<double> const& a7,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(283): error: complex is not a template
        complex<double> const& b0, complex<double> const& b1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(283): error: complex is not a template
        complex<double> const& b0, complex<double> const& b1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(284): error: complex is not a template
        complex<double> const& b2, complex<double> const& b3,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(284): error: complex is not a template
        complex<double> const& b2, complex<double> const& b3,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(285): error: complex is not a template
        complex<double> const& c0, complex<double> const& c1,
        ^

cutlass/include/cute/arch/mma_sm90.hpp(285): error: complex is not a template
        complex<double> const& c0, complex<double> const& c1,
                                   ^

cutlass/include/cute/arch/mma_sm90.hpp(286): error: complex is not a template
        complex<double> const& c2, complex<double> const& c3)
        ^

cutlass/include/cute/arch/mma_sm90.hpp(286): error: complex is not a template
        complex<double> const& c2, complex<double> const& c3)
                                   ^

cutlass/include/cutlass/epilogue/fusion/operations.hpp(191): error: expected a ","
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                                         ^
          detected during processing of template argument list for "cutlass::epilogue::fusion::LinCombPerRowBiasEltActAux" based on template arguments <GmemLayoutTagD, Schedule::ActivationFunctor, ElementD, ElementCompute, Schedule::ElementT, Schedule::ElementBias, ElementC_, ElementCompute, <error-constant>> at line 696 of cutlass/include/cutlass/epilogue/collective/builders/sm90_builder.inl

cutlass/include/cutlass/epilogue/fusion/operations.hpp(152): error: type name is not allowed
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                             ^
          detected during processing of template argument list for "cutlass::epilogue::fusion::LinCombPerRowBiasEltAct" based on template arguments <Schedule::ActivationFunctor, ElementD, ElementCompute, Schedule::ElementBias, ElementC_, ElementCompute> at line 704 of cutlass/include/cutlass/epilogue/collective/builders/sm90_builder.inl

cutlass/include/cutlass/epilogue/fusion/operations.hpp(152): error: expected a ","
    int AlignmentBias_ = 128 / sizeof_bits_v<ElementBias_>,
                                                         ^
          detected during processing of template argument list for "cutlass::epilogue::fusion::LinCombPerRowBiasEltAct" based on template arguments <Schedule::ActivationFunctor, ElementD, ElementCompute, Schedule::ElementBias, ElementC_, ElementCompute> at line 704 of cutlass/include/cutlass/epilogue/collective/builders/sm90_builder.inl

83 errors detected in the compilation of "test.cu".

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with the standalone include example, cutlass/epilogue/collective/collective_builder.hpp, and reproduce it with nvcc using the listed include paths and C++20. Review the shell-script-generated list and verify each header can be included individually; done means the listed headers compile cleanly without relying on transitive includes.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
build-system, testing-qa
Issue type
Bug
Difficulty
5/5
Estimated time
Over a week
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.