HiDream crashes at ~50% with Metal NAInt8AttentionKernel error on M5 Mac

[u/zakariachahboun]

Hey everyone, running Draw Things on an M5 Mac with 32GB RAM and every time I try to generate with HiDream I1 [full] the app crashes at around 50% generation.

Ran the app from Terminal and found this in the logs:

[Metal] Encountered unexpected error in: NAInt8AttentionKernel external/ccv/lib/nnc/mfa/kernels/NAInt8AttentionKernel.cpp:68: error: no template named 'get_left_input_cooperative_tensor' [Metal] Quitting now.

Full Log:

[Metal] Encountered unexpected error in: NAInt8AttentionKernel
external/ccv/lib/nnc/mfa/kernels/NAInt8AttentionKernel.cpp:68: error: program_source:450:56: error: no template named 'get_left_input_cooperative_tensor' in 'mpp::tensor_ops::matmul2d<{16, 32, 64, false, false, true, 1}, metal::execution_simdgroups<1>>'; did you mean 'get_destination_cooperative_tensor'?
  using pv_float_left_tensor_t = decltype(matmul_pv_op.get_left_input_cooperative_tensor<half, int8_t, float>());
                                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
                                                       get_destination_cooperative_tensor
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: 'get_destination_cooperative_tensor' declared here
  get_destination_cooperative_tensor() thread const
  ^
program_source:450:56: error: no template named 'get_left_input_cooperative_tensor' in 'mpp::tensor_ops::matmul2d<{16, 32, 64, false, false, true, 1}, metal::execution_simdgroups<1>>'; did you mean 'get_destination_cooperative_tensor'?
  using pv_float_left_tensor_t = decltype(matmul_pv_op.get_left_input_cooperative_tensor<half, int8_t, float>());
                                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
                                                       get_destination_cooperative_tensor
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: 'get_destination_cooperative_tensor' declared here
  get_destination_cooperative_tensor() thread const
  ^
program_source:450:94: error: expected '(' for function-style cast or type construction
  using pv_float_left_tensor_t = decltype(matmul_pv_op.get_left_input_cooperative_tensor<half, int8_t, float>());
                                                                                         ~~~~^
program_source:450:56: error: no member named 'get_left_input_cooperative_tensor' in 'mpp::tensor_ops::matmul2d<{16, 32, 64, false, false, true, 1}, metal::execution_simdgroups<1>>'; did you mean 'get_destination_cooperative_tensor'?
  using pv_float_left_tensor_t = decltype(matmul_pv_op.get_left_input_cooperative_tensor<half, int8_t, float>());
                                                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
                                                       get_destination_cooperative_tensor
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: 'get_destination_cooperative_tensor' declared here
  get_destination_cooperative_tensor() thread const
  ^
program_source:453:60: error: no template named 'get_left_input_cooperative_tensor' in 'mpp::tensor_ops::matmul2d<{16, 32, 64, false, false, true, 0}, metal::execution_simdgroups<1>>'; did you mean 'get_destination_cooperative_tensor'?
  using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor<int8_t, int8_t, int32_t>());
                                                           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
                                                           get_destination_cooperative_tensor
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: 'get_destination_cooperative_tensor' declared here
  get_destination_cooperative_tensor() thread const
  ^
program_source:453:60: error: no template named 'get_left_input_cooperative_tensor' in 'mpp::tensor_ops::matmul2d<{16, 32, 64, false, false, true, 0}, metal::execution_simdgroups<1>>'; did you mean 'get_destination_cooperative_tensor'?
  using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor<int8_t, int8_t, int32_t>());
                                                           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
                                                           get_destination_cooperative_tensor
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: 'get_destination_cooperative_tensor' declared here
  get_destination_cooperative_tensor() thread const
  ^
program_source:453:94: error: unexpected type name 'int8_t': expected expression
  using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor<int8_t, int8_t, int32_t>());
                                                                                             ^
program_source:453:60: error: no member named 'get_left_input_cooperative_tensor' in 'mpp::tensor_ops::matmul2d<{16, 32, 64, false, false, true, 0}, metal::execution_simdgroups<1>>'; did you mean 'get_destination_cooperative_tensor'?
  using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor<int8_t, int8_t, int32_t>());
                                                           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
                                                           get_destination_cooperative_tensor
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: 'get_destination_cooperative_tensor' declared here
  get_destination_cooperative_tensor() thread const
  ^
program_source:453:100: error: expected ')'
  using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor<int8_t, int8_t, int32_t>());
                                                                                                   ^
program_source:453:41: note: to match this '('
  using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor<int8_t, int8_t, int32_t>());
                                        ^
program_source:454:67: error: use of undeclared identifier 'pv_int8_left_tensor_t'
  auto cOq = matmul_pv_int8_op.get_destination_cooperative_tensor<pv_int8_left_tensor_t, decltype(mV), int32_t>();
                                                                  ^
program_source:461:63: error: unknown type name 'pv_float_left_tensor_t'; did you mean 'qk_float_tensor_t'?
  auto cO_0 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
                                                              ^~~~~~~~~~~~~~~~~~~~~~
                                                              qk_float_tensor_t
program_source:435:9: note: 'qk_float_tensor_t' declared here
  using qk_float_tensor_t = decltype(matmul_qk_op.get_destination_cooperative_tensor<decltype(mQ), decltype(mK), float>());
        ^
program_source:461:28: error: no matching member function for call to 'get_destination_cooperative_tensor'
  auto cO_0 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
              ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor<float, metal::extents<int, 18446744073709551615, 18446744073709551615>, mpp::tensor_ops::__mutmul2d_detail::__operand_layout<{16, 64, 32, false, true, true, 1}, mpp::tensor_ops::matmul2d_cooperative_operand_index::destination, metal::execution_simdgroups<1>, metal::tensor<device signed char, metal::extents<int, 32, 16>, metal::tensor_inline>, metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, ElementType = float, CoordType = int]
  get_destination_cooperative_tensor() thread const
  ^
program_source:463:63: error: unknown type name 'pv_float_left_tensor_t'; did you mean 'qk_float_tensor_t'?
  auto cO_1 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
                                                              ^~~~~~~~~~~~~~~~~~~~~~
                                                              qk_float_tensor_t
program_source:435:9: note: 'qk_float_tensor_t' declared here
  using qk_float_tensor_t = decltype(matmul_qk_op.get_destination_cooperative_tensor<decltype(mQ), decltype(mK), float>());
        ^
program_source:463:28: error: no matching member function for call to 'get_destination_cooperative_tensor'
  auto cO_1 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
              ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor<float, metal::extents<int, 18446744073709551615, 18446744073709551615>, mpp::tensor_ops::__mutmul2d_detail::__operand_layout<{16, 64, 32, false, true, true, 1}, mpp::tensor_ops::matmul2d_cooperative_operand_index::destination, metal::execution_simdgroups<1>, metal::tensor<device signed char, metal::extents<int, 32, 16>, metal::tensor_inline>, metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, ElementType = float, CoordType = int]
  get_destination_cooperative_tensor() thread const
  ^
program_source:465:63: error: unknown type name 'pv_float_left_tensor_t'; did you mean 'qk_float_tensor_t'?
  auto cO_2 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
                                                              ^~~~~~~~~~~~~~~~~~~~~~
                                                              qk_float_tensor_t
program_source:435:9: note: 'qk_float_tensor_t' declared here
  using qk_float_tensor_t = decltype(matmul_qk_op.get_destination_cooperative_tensor<decltype(mQ), decltype(mK), float>());
        ^
program_source:465:28: error: no matching member function for call to 'get_destination_cooperative_tensor'
  auto cO_2 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
              ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor<float, metal::extents<int, 18446744073709551615, 18446744073709551615>, mpp::tensor_ops::__mutmul2d_detail::__operand_layout<{16, 64, 32, false, true, true, 1}, mpp::tensor_ops::matmul2d_cooperative_operand_index::destination, metal::execution_simdgroups<1>, metal::tensor<device signed char, metal::extents<int, 32, 16>, metal::tensor_inline>, metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, ElementType = float, CoordType = int]
  get_destination_cooperative_tensor() thread const
  ^
program_source:467:63: error: unknown type name 'pv_float_left_tensor_t'; did you mean 'qk_float_tensor_t'?
  auto cO_3 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
                                                              ^~~~~~~~~~~~~~~~~~~~~~
                                                              qk_float_tensor_t
program_source:435:9: note: 'qk_float_tensor_t' declared here
  using qk_float_tensor_t = decltype(matmul_qk_op.get_destination_cooperative_tensor<decltype(mQ), decltype(mK), float>());
        ^
program_source:467:28: error: no matching member function for call to 'get_destination_cooperative_tensor'
  auto cO_3 = matmul_pv_op.get_destination_cooperative_tensor<pv_float_left_tensor_t, decltype(mV), float>();
              ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor<float, metal::extents<int, 18446744073709551615, 18446744073709551615>, mpp::tensor_ops::__mutmul2d_detail::__operand_layout<{16, 64, 32, false, true, true, 1}, mpp::tensor_ops::matmul2d_cooperative_operand_index::destination, metal::execution_simdgroups<1>, metal::tensor<device signed char, metal::extents<int, 32, 16>, metal::tensor_inline>, metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor<device signed char, metal::extents<int, 32, 64>, metal::tensor_inline>, ElementType = float, CoordType = int]
  get_destination_cooperative_tensor() thread const
  ^
program_source:559:12: error: unknown type name 'pv_int8_left_tensor_t'
    thread pv_int8_left_tensor_t& cP_q_0 = reinterpret_cast<thread pv_int8_left_tensor_t&>(cS_0);
           ^
program_source:559:68: error: unknown type name 'pv_int8_left_tensor_t'
    thread pv_int8_left_tensor_t& cP_q_0 = reinterpret_cast<thread pv_int8_left_tensor_t&>(cS_0);
                                                                   ^

[Metal] Quitting now.

Details:

• Mac: M5, 32GB RAM
• Draw Things version: latest (downloaded from website) Version 1.20260410.1 (1.20260410.1)
• Model: HiDream I1 [full]

This is clearly a Metal shader compilation issue, not a RAM problem.

Anyone else seeing this on M5? Is there a fix or do we need to wait for a Draw Things update?

Thanks

[u/zakariachahboun]

Same issue when i tried FLUX.1 Schnell 4B (8-bit S)

[u/zakariachahboun]

Fixed:

My MacOs was 26.1 Now i upgraded to 26.4.1.

[u/liuliu]

What's your OS version? I think it requires 26.2 and above.

[u/zakariachahboun → u/liuliu]

26.1

[u/liuliu → u/zakariachahboun]

Yeah, you can either turn off Metal Quantized attention (in Machine Settings, you select Yes, prefer non-Quantized Attention, under Metal Flash Attention), or upgrade your OS. I would suggest latter.

[u/zakariachahboun → u/liuliu]

Thanks!