Crashes with Metal NAInt8AttentionKernel error on M5 Mac

[zakaria-chahboun]

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()); ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 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()); ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 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()); ~~~~^ 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()); ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 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()); ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 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()); ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 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()); ^ 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()); ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 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()); ^ program_source:453:41: note: to match this '(' using pv_int8_left_tensor_t = decltype(matmul_pv_int8_op.get_left_input_cooperative_tensor()); ^ program_source:454:67: error: use of undeclared identifier 'pv_int8_left_tensor_t' auto cOq = matmul_pv_int8_op.get_destination_cooperative_tensor(); ^ 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(); ^~~~~~~~~~~~~~~~~~~~~~ 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()); ^ 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(); ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ /System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor, 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, metal::tensor_inline>, metal::tensor, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor, 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(); ^~~~~~~~~~~~~~~~~~~~~~ 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()); ^ 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(); ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ /System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor, 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, metal::tensor_inline>, metal::tensor, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor, 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(); ^~~~~~~~~~~~~~~~~~~~~~ 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()); ^ 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(); ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ /System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor, 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, metal::tensor_inline>, metal::tensor, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor, 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(); ^~~~~~~~~~~~~~~~~~~~~~ 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()); ^ 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(); ~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ /System/Library/Frameworks/MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h:421:3: note: candidate template ignored: requirement 'false' was not satisfied [with LeftOperandType = metal::cooperative_tensor, 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, metal::tensor_inline>, metal::tensor, metal::tensor_inline>, float, int>>, RightOperandType = metal::tensor, 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(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(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

[zakaria-chahboun]

Same things when i tried another model like FLUX.1 Schnell 4B (8-bit S)

[zakaria-chahboun]

Fixed:

My MacOs was 26.1 Now i upgraded to 26.4.1.