code: def test_translate_ethosu_depthwise_conv2d():
      def extract_ethosu_depthwise_conv2d_extern_call(mod):
          # There should only be a single function
          assert len(mod.functions.items()) == 1
          primfunc = mod.functions.items()[0][1]
  
          ethosu_depthwise_conv2d_calls = list()
  
          def populate_ethosu_depthwise_conv2d_calls(stmt):
              if (
                  isinstance(stmt, tvm.tir.Call)
                  and stmt.op.name == "tir.call_extern"
                  and stmt.args[0] == "ethosu_depthwise_conv2d"
              ):
                  ethosu_depthwise_conv2d_calls.append(stmt)
  
          stmt_functor.post_order_visit(
              primfunc.body, populate_ethosu_depthwise_conv2d_calls
          )
          return ethosu_depthwise_conv2d_calls[0]
  
      depthwise_conv2d_call = (
          extract_ethosu_depthwise_conv2d_extern_call(
              SingleEthosuDepthwiseConv2D
          )
      )
      (
          npu_op,
          w_zero_point,
      ) = tir_to_cs_translator.translate_ethosu_depthwise_conv2d(
          depthwise_conv2d_call
      )
  
      assert npu_op.ifm.data_type == vapi.NpuDataType.INT8
      assert npu_op.ifm.shape == vapi.NpuShape3D(8, 8, 3)
      assert (
          npu_op.ifm.tiles.height_0
          == vapi.NpuTileBox(8, 0, 8, [0, 0, 0, 0]).height_0
      )
      assert (
          npu_op.ifm.tiles.height_1
          == vapi.NpuTileBox(8, 0, 8, [0, 0, 0, 0]).height_1
      )
      assert (
          npu_op.ifm.tiles.width_0
          == vapi.NpuTileBox(8, 0, 8, [0, 0, 0, 0]).width_0
      )
      assert npu_op.ifm.quantization == pytest.approx(
          vapi.NpuQuantization(0.6, 11)
      )
      assert npu_op.ifm.layout == vapi.NpuLayout.NHWC
      assert npu_op.ifm.strides == vapi.NpuShape3D(24, 3, 1)
      # Compare OFM
      assert npu_op.ofm.data_type == vapi.NpuDataType.INT8
      assert npu_op.ofm.shape == vapi.NpuShape3D(6, 7, 3)
      assert (
          npu_op.ofm.tiles.height_0
          == vapi.NpuTileBox(6, 0, 8, [0, 0, 0, 0]).height_0
      )
      assert (
          npu_op.ofm.tiles.height_1
          == vapi.NpuTileBox(6, 0, 7, [0, 0, 0, 0]).height_1
      )
      assert (
          npu_op.ofm.tiles.width_0
          == vapi.NpuTileBox(6, 0, 7, [0, 0, 0, 0]).width_0
      )
      assert npu_op.ofm.quantization == pytest.approx(
          vapi.NpuQuantization(0.26, 15)
      )
      assert npu_op.ofm.layout == vapi.NpuLayout.NHWC
      assert npu_op.ofm.strides == vapi.NpuShape3D(21, 3, 1)
      # Compare kernel and padding
      assert (
          npu_op.kernel.__dict__
          == vapi.NpuKernel(
              w=2,
              h=3,
              stride_x=1,
              stride_y=1,
              dilation_x=1,
              dilation_y=1,
          ).__dict__
      )
      assert npu_op.padding == vapi.NpuPadding(
          top=0, left=0, bottom=0, right=0
      )
      # Compare activation
      assert (
          npu_op.activation.op_type == vapi.NpuActivationOp.NONE_OR_RELU
      )
      assert npu_op.activation.min == 0
      assert npu_op.activation.max == pytest.approx(23.4)
      # Compare rounding mode
      assert npu_op.rounding_mode == vapi.NpuRoundingMode.TFL
      # Compare ifm upscaling
      assert npu_op.ifm_upscale == vapi.NpuResamplingMode.NONE
      # Compare weight quantization parameters
      assert w_zero_point == 13
  
imports: import pytest
  import numpy
  import tvm
  from tvm.tir import stmt_functor
  from tvm.script import tir
  from tvm.relay.backend.contrib.ethosu import tir_to_cs_translator
  from tvm.relay.backend.contrib.ethosu import util
  import ethosu.vela.api
label_desc: Use image node_id 3 for calling a global function with example usage: test_translate_ethosu_depthwise_conv2d() and returns: ethosu_depthwise_conv2d_calls