Repository: llvm/llvm-project
File: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (C++)

Code before:
  case Intrinsic::nvvm_tcgen05_st_16x128b_x16:
  case Intrinsic::nvvm_tcgen05_st_16x128b_x32:
  case Intrinsic::nvvm_tcgen05_st_16x128b_x64:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x1:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x2:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x4:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x32:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x16:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x32:
  case Intrinsic::nvvm_tcgen05_st_16x64b_x64:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
    return lowerTcgen05St(Op, DAG);
  case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
  case Intrinsic::
      nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
  case Intrinsic::
      nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:

Review comment (suggestion):
so, these were lowered with hasOffset=false previously?
If we need new tests for this update, please add them.

Change itself LGTM

Code after:
  case Intrinsic::nvvm_tcgen05_st_16x128b_x16:
  case Intrinsic::nvvm_tcgen05_st_16x128b_x32:
  case Intrinsic::nvvm_tcgen05_st_16x128b_x64:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x1:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x2:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x4:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
  case Intrinsic::nvvm_tcgen05_st_16x256b_x32:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x16:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x32:
  case Intrinsic::nvvm_tcgen05_st_16x64b_x64:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
  case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
    return lowerTcgen05St(Op, DAG);
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64:
  case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128:
    return lowerTcgen05St(Op, DAG, /* hasOffset */ true);
  case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
  case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
  case Intrinsic::
      nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
  case Intrinsic::
      nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
  case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
  case Intrinsic::