code: def measure_bandwidth_sum(
      total_item,
      item_per_thread,
      stride,
      base_type,
      bits,
      lanes,
      target,
      target_host,
      remote,
      dev,
      n_times,
  ):
      """measure memory bandwidth of gpu by product reduction for a given type
  
      The IR for measurement is
  
      for each thread
          for i in 1..num_per_thread:
              y[global_id] = y[global_id] * x[base + i * stride]
  
      Parameters
      ----------
      total_item: int
          number of elements in input array
      item_per_thread: int
          number of elements each thread accumulates
      stride: int
          stride in memory access
      base_type: str
          can be "int", "float"
      bits: int
          can be 16, 32
      lanes: int
         lane of the vector type, can be 1, 2, 4, 8, 16
      target: :any:`tvm.target.Target`
          the target and option of the compilation.
      target_host : str or :any:`tvm.target.Target`
          host compilation target
      dev: Device
          the device of array
      remote: tvm.rpc.RPCSession
          remote rpc session
      n_times: int
          number of runs for taking mean
  
      Returns
      -------
      GBPS: float
           gigabyte per second
      """
      target, target_host = Target.canon_target_and_host(
          target, target_host
      )
  
      n, m = total_item, item_per_thread
      n //= lanes
  
      base_type = str(base_type) + str(bits)
      dtype = base_type if lanes == 1 else base_type + "x" + str(lanes)
  
      k = te.reduce_axis((0, m), name="k")
  
      x = te.placeholder((n,), dtype=dtype, name="x")
      op = te.comm_reducer(
          lambda x, y: x * y,
          lambda t: tvm.tir.const(1, dtype=t),
          name="sum",
      )
      y = te.compute(
          (n // m,),
          lambda i: op(
              x[i // stride * stride * m + i % stride + k * stride],
              axis=k,
          ),
      )
      s = te.create_schedule(y.op)
  
      yo, yi = s[y].split(y.op.axis[0], target.max_num_threads)
      s[y].bind(yo, te.thread_axis("blockIdx.x"))
      s[y].bind(yi, te.thread_axis("threadIdx.x"))
      s[y].unroll(k)
  
      try:
          func = tvm.build(s, [x, y], target)
  
          x = tvm.nd.empty((n,), dtype=dtype, device=dev)
          y = tvm.nd.empty((n // m,), dtype=dtype, device=dev)
  
          func = _convert_to_remote(func, remote)
          time_f = func.time_evaluator(
              func.entry_name, dev, number=n_times
          )
          time = time_f(x, y).mean
      except tvm._ffi.base.TVMError:
          # build error (occur when device does not support half)
          return -1
  
      return 1.0 * (total_item * bits / 8) / 1e9 / time
  
file_path: tvm/python/tvm/contrib/peak.py
imports: import logging
  import tvm
  from tvm import te
  from tvm.target import Target
  from .None import utils
  from ..None import rpc
label_desc: Use image node_id 2 for calling a global function with example usage: measure_bandwidth_sum(total_item, item_per_thread, stride, base_type, bits, lanes, target, target_host, remote, dev, n_times) and returns: unknown, unknown