code: def index_select_cat_bwd_kernel(
      grad_source_ptr,  # *Pointer* to grad_source tensor.
      index_ptr,  # *Pointer* to index tensor.
      grad_output_ptr,  # *Pointer* to grad_output tensor.
      num_rows,
      num_indices,
      num_cols,
      stride0,  # Stride information of input and source tensor.
      stride1,
      BLOCK_SIZE_INDEX: tl.constexpr,  # Number of indices each program should process.
      BLOCK_SIZE_COL: tl.constexpr,  # Number of cols each program should process.
  ):
      pid0 = tl.program_id(axis=0)  # We use 3D launch grid
      pid1 = tl.program_id(axis=1)
  
      cols = pid1 * BLOCK_SIZE_COL + tl.arange(0, BLOCK_SIZE_COL)
  
      # load grad_output
      grad_output_indices = pid0 * BLOCK_SIZE_INDEX + tl.arange(
          0, BLOCK_SIZE_INDEX
      )
      grad_output_offsets = (
          grad_output_ptr
          + grad_output_indices[:, None] * stride0
          + cols[None, :] * stride1
      )
      grad_output_mask = (
          grad_output_indices[:, None] < num_indices
      ) & (cols[None, :] < num_cols)
      grad_output = tl.load(
          grad_output_offsets, mask=grad_output_mask
      ).to(tl.float32)
  
      # select indices from grad_source
      grad_source_indices = tl.load(
          index_ptr + grad_output_indices,
          mask=(grad_output_indices < num_indices),
      )
      grad_source_offsets = (
          grad_source_ptr
          + grad_source_indices[:, None] * stride0
          + cols[None, :] * stride1
      )
  
      # compute scaled index add and save
      tl.store(grad_source_offsets, grad_output, mask=grad_output_mask)
  
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 3 for calling a global function with example usage: index_select_cat_bwd_kernel(grad_source_ptr, index_ptr, grad_output_ptr, num_rows, num_indices, num_cols, stride0, stride1, BLOCK_SIZE_INDEX, BLOCK_SIZE_COL) without return types