code: def l2_step1_k(indptr, data, norm):
      """Compute norm for L2 normalization"""
      row_i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x
      inrow_idx = cuda.blockIdx.y * cuda.blockDim.y + cuda.threadIdx.y
  
      if row_i >= indptr.shape[0] - 1:
          return
  
      start = indptr[row_i]
      end = indptr[row_i + 1]
      if inrow_idx >= (end - start):
          return
  
      val = data[start + inrow_idx]
      val *= val
      cuda.atomic.add(norm, row_i, val)
  
file_path: cuml/python/cuml/thirdparty_adapters/sparsefuncs_fast.py
imports: from math import ceil
  from cuml.internals.safe_imports import gpu_only_import_from
  from cuml.internals.safe_imports import gpu_only_import
label_desc: Use image node_id 7 for calling a global function with example usage: l2_step1_k(indptr, data, norm) without return types