content: import triton
  import triton.language as tl
  from triton.compiler.compiler import AttrsDescriptor
  
  from torch._inductor.runtime import triton_helpers, triton_heuristics
  from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
  from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, DeviceProperties
  
  @triton_heuristics.template(
      num_stages=2,
      num_warps=1,
      triton_meta={'signature': {'arg_A': '*fp32', 'arg_B': '*fp32', 'in_ptr2': '*fp32', 'in_ptr3': '*fp32', 'out_ptr1': '*fp32'}, 'device': DeviceProperties(type='cuda', index=0, multi_processor_count=78, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, warp_size=32), 'constants': {}, 'configs': [AttrsDescriptor.from_dict({'arg_properties': {'tt.divisibility': (0, 1, 2, 3, 4), 'tt.equal_to': ()}, 'cls': 'AttrsDescriptor'})]},
      inductor_meta={'kernel_name': 'triton_tem_fused_add_addmm_mm_mul_sigmoid_5', 'backend_hash': '9DCB932EF217650ABAB636BFB7B75C0464066E10820E4CB71C69D98802C3B71B', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': True, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False, 'coordinate_descent_tuning': True, 'coordinate_descent_search_radius': 1, 'coordinate_descent_check_all_directions': False},
  )
  @triton.jit
  def triton_tem_fused_add_addmm_mm_mul_sigmoid_5(arg_A, arg_B, in_ptr2, in_ptr3, out_ptr1):
      GROUP_M : tl.constexpr = 8
      EVEN_K : tl.constexpr = False
      ALLOW_TF32 : tl.constexpr = False
      ACC_TYPE : tl.constexpr = tl.float32
      B_PROLOGUE_CAST_TYPE : tl.constexpr = None
      BLOCK_M : tl.constexpr = 16
      BLOCK_N : tl.constexpr = 16
      BLOCK_K : tl.constexpr = 16
      A = arg_A
      B = arg_B
  
      M = 7
      N = 4
      K = 5
      if M * N == 0:
          # early exit due to zero-size input(s)
          return
      stride_am = 5
      stride_ak = 1
      stride_bk = 4
      stride_bn = 1
  
      # based on triton.ops.matmul
      pid = tl.program_id(0)
      grid_m = (M + BLOCK_M - 1) // BLOCK_M
      grid_n = (N + BLOCK_N - 1) // BLOCK_N
  
      # re-order program ID for better L2 performance
      width = GROUP_M * grid_n
      group_id = pid // width
      group_size = min(grid_m - group_id * GROUP_M, GROUP_M)
      pid_m = group_id * GROUP_M + (pid % group_size)
      pid_n = (pid % width) // (group_size)
  
      rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
      rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
      if (stride_am == 1 and stride_ak == M) or (stride_am == K and stride_ak == 1):
          ram = tl.max_contiguous(tl.multiple_of(rm % M, BLOCK_M), BLOCK_M)
      else:
          ram = rm % M
      if (stride_bk == 1 and stride_bn == K) or (stride_bk == N and stride_bn == 1):
          rbn = tl.max_contiguous(tl.multiple_of(rn % N, BLOCK_N), BLOCK_N)
      else:
          rbn = rn % N
      rk = tl.arange(0, BLOCK_K)
      A = A + (ram[:, None] * stride_am + rk[None, :] * stride_ak)
      B = B + (rk[:, None] * stride_bk + rbn[None, :] * stride_bn)
  
      acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=ACC_TYPE)
      for k in range(K, 0, -BLOCK_K):
          if EVEN_K:
              a = tl.load(A)
              b = tl.load(B)
          else:
              a = tl.load(A, mask=rk[None, :] < k, other=0.)
              b = tl.load(B, mask=rk[:, None] < k, other=0.)
          if B_PROLOGUE_CAST_TYPE is not None:
              b = b.to(B_PROLOGUE_CAST_TYPE)
          acc += tl.dot(a, b, allow_tf32=ALLOW_TF32)
          A += BLOCK_K * stride_ak
          B += BLOCK_K * stride_bk
  
      # rematerialize rm and rn to save registers
      rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
      rn = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
      idx_m = rm[:, None]
      idx_n = rn[None, :]
      mask = (idx_m < M) & (idx_n < N)
  
      # inductor generates a suffix
      xindex = idx_n + 4*idx_m
      tmp0 = tl.load(in_ptr2 + (tl.broadcast_to(idx_n + 4*idx_m, acc.shape)), mask, eviction_policy='evict_last')
      tmp1 = tl.load(in_ptr3 + (tl.broadcast_to(idx_n, acc.shape)), mask, eviction_policy='evict_last')
      tmp2 = tmp0 + tmp1
      tmp3 = 0.5
      tmp4 = acc * tmp3
      tmp5 = tmp2 + tmp4
      tmp6 = tl.sigmoid(tmp5)
      tl.store(out_ptr1 + (tl.broadcast_to(idx_n + 4*idx_m, acc.shape)), tmp6, mask)