in llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp [3436:4281]
bool NVPTXTargetLowering::getTgtMemIntrinsic(
IntrinsicInfo &Info, const CallInst &I,
MachineFunction &MF, unsigned Intrinsic) const {
switch (Intrinsic) {
default:
return false;
case Intrinsic::nvvm_match_all_sync_i32p:
case Intrinsic::nvvm_match_all_sync_i64p:
Info.opc = ISD::INTRINSIC_W_CHAIN;
// memVT is bogus. These intrinsics have IntrInaccessibleMemOnly attribute
// in order to model data exchange with other threads, but perform no real
// memory accesses.
Info.memVT = MVT::i1;
// Our result depends on both our and other thread's arguments.
Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore;
return true;
case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v8f16;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v2i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(8);
return true;
}
case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row_stride:
case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col:
case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row:
case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row_stride:
case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col:
case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row:
case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row_stride:
case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16:
case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v4i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row:
case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row:
case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row_stride:
case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col:
case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row:
case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row:
case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col:
case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col:
case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16:
case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(4);
return true;
}
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v4f16;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row_stride:
case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col:
case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row:
case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v8f32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row:
case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row:
case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row:
case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v8i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col:
case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col_stride:
case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row:
case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col:
case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col_stride:
case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row:
case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row_stride:
case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16:
case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v2i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(8);
return true;
}
case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col:
case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col_stride:
case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row:
case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row_stride:
case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col:
case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col_stride:
case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row:
case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::f64;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(8);
return true;
}
case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col:
case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col_stride:
case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row:
case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row_stride: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v2f64;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row_stride: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v4f16;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOStore;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row_stride:
case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col:
case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row:
case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row_stride: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v8f32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOStore;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col_stride:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row:
case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row_stride:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col_stride:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row:
case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row_stride:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col_stride:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row:
case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row_stride: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v8i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOStore;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col:
case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col_stride:
case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row:
case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row_stride:
case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col:
case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col_stride:
case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row:
case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row_stride: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v2i32;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOStore;
Info.align = Align(8);
return true;
}
case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col:
case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col_stride:
case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row:
case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row_stride: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v2f64;
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOStore;
Info.align = Align(16);
return true;
}
case Intrinsic::nvvm_atomic_load_inc_32:
case Intrinsic::nvvm_atomic_load_dec_32:
case Intrinsic::nvvm_atomic_add_gen_f_cta:
case Intrinsic::nvvm_atomic_add_gen_f_sys:
case Intrinsic::nvvm_atomic_add_gen_i_cta:
case Intrinsic::nvvm_atomic_add_gen_i_sys:
case Intrinsic::nvvm_atomic_and_gen_i_cta:
case Intrinsic::nvvm_atomic_and_gen_i_sys:
case Intrinsic::nvvm_atomic_cas_gen_i_cta:
case Intrinsic::nvvm_atomic_cas_gen_i_sys:
case Intrinsic::nvvm_atomic_dec_gen_i_cta:
case Intrinsic::nvvm_atomic_dec_gen_i_sys:
case Intrinsic::nvvm_atomic_inc_gen_i_cta:
case Intrinsic::nvvm_atomic_inc_gen_i_sys:
case Intrinsic::nvvm_atomic_max_gen_i_cta:
case Intrinsic::nvvm_atomic_max_gen_i_sys:
case Intrinsic::nvvm_atomic_min_gen_i_cta:
case Intrinsic::nvvm_atomic_min_gen_i_sys:
case Intrinsic::nvvm_atomic_or_gen_i_cta:
case Intrinsic::nvvm_atomic_or_gen_i_sys:
case Intrinsic::nvvm_atomic_exch_gen_i_cta:
case Intrinsic::nvvm_atomic_exch_gen_i_sys:
case Intrinsic::nvvm_atomic_xor_gen_i_cta:
case Intrinsic::nvvm_atomic_xor_gen_i_sys: {
auto &DL = I.getModule()->getDataLayout();
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = getValueType(DL, I.getType());
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad | MachineMemOperand::MOStore;
Info.align.reset();
return true;
}
case Intrinsic::nvvm_ldu_global_i:
case Intrinsic::nvvm_ldu_global_f:
case Intrinsic::nvvm_ldu_global_p: {
auto &DL = I.getModule()->getDataLayout();
Info.opc = ISD::INTRINSIC_W_CHAIN;
if (Intrinsic == Intrinsic::nvvm_ldu_global_i)
Info.memVT = getValueType(DL, I.getType());
else if(Intrinsic == Intrinsic::nvvm_ldu_global_p)
Info.memVT = getPointerTy(DL);
else
Info.memVT = getValueType(DL, I.getType());
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = cast<ConstantInt>(I.getArgOperand(1))->getMaybeAlignValue();
return true;
}
case Intrinsic::nvvm_ldg_global_i:
case Intrinsic::nvvm_ldg_global_f:
case Intrinsic::nvvm_ldg_global_p: {
auto &DL = I.getModule()->getDataLayout();
Info.opc = ISD::INTRINSIC_W_CHAIN;
if (Intrinsic == Intrinsic::nvvm_ldg_global_i)
Info.memVT = getValueType(DL, I.getType());
else if(Intrinsic == Intrinsic::nvvm_ldg_global_p)
Info.memVT = getPointerTy(DL);
else
Info.memVT = getValueType(DL, I.getType());
Info.ptrVal = I.getArgOperand(0);
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = cast<ConstantInt>(I.getArgOperand(1))->getMaybeAlignValue();
return true;
}
case Intrinsic::nvvm_tex_1d_v4f32_s32:
case Intrinsic::nvvm_tex_1d_v4f32_f32:
case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
case Intrinsic::nvvm_tex_2d_v4f32_s32:
case Intrinsic::nvvm_tex_2d_v4f32_f32:
case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
case Intrinsic::nvvm_tex_3d_v4f32_s32:
case Intrinsic::nvvm_tex_3d_v4f32_f32:
case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
case Intrinsic::nvvm_tex_cube_v4f32_f32:
case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
Info.opc = getOpcForTextureInstr(Intrinsic);
Info.memVT = MVT::v4f32;
Info.ptrVal = nullptr;
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
case Intrinsic::nvvm_tex_1d_v4s32_s32:
case Intrinsic::nvvm_tex_1d_v4s32_f32:
case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
case Intrinsic::nvvm_tex_2d_v4s32_s32:
case Intrinsic::nvvm_tex_2d_v4s32_f32:
case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
case Intrinsic::nvvm_tex_3d_v4s32_s32:
case Intrinsic::nvvm_tex_3d_v4s32_f32:
case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
case Intrinsic::nvvm_tex_cube_v4s32_f32:
case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
case Intrinsic::nvvm_tex_cube_v4u32_f32:
case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
case Intrinsic::nvvm_tex_1d_v4u32_s32:
case Intrinsic::nvvm_tex_1d_v4u32_f32:
case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
case Intrinsic::nvvm_tex_2d_v4u32_s32:
case Intrinsic::nvvm_tex_2d_v4u32_f32:
case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
case Intrinsic::nvvm_tex_3d_v4u32_s32:
case Intrinsic::nvvm_tex_3d_v4u32_f32:
case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
Info.opc = getOpcForTextureInstr(Intrinsic);
Info.memVT = MVT::v4i32;
Info.ptrVal = nullptr;
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
case Intrinsic::nvvm_suld_1d_i8_clamp:
case Intrinsic::nvvm_suld_1d_v2i8_clamp:
case Intrinsic::nvvm_suld_1d_v4i8_clamp:
case Intrinsic::nvvm_suld_1d_array_i8_clamp:
case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
case Intrinsic::nvvm_suld_2d_i8_clamp:
case Intrinsic::nvvm_suld_2d_v2i8_clamp:
case Intrinsic::nvvm_suld_2d_v4i8_clamp:
case Intrinsic::nvvm_suld_2d_array_i8_clamp:
case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
case Intrinsic::nvvm_suld_3d_i8_clamp:
case Intrinsic::nvvm_suld_3d_v2i8_clamp:
case Intrinsic::nvvm_suld_3d_v4i8_clamp:
case Intrinsic::nvvm_suld_1d_i8_trap:
case Intrinsic::nvvm_suld_1d_v2i8_trap:
case Intrinsic::nvvm_suld_1d_v4i8_trap:
case Intrinsic::nvvm_suld_1d_array_i8_trap:
case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
case Intrinsic::nvvm_suld_2d_i8_trap:
case Intrinsic::nvvm_suld_2d_v2i8_trap:
case Intrinsic::nvvm_suld_2d_v4i8_trap:
case Intrinsic::nvvm_suld_2d_array_i8_trap:
case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
case Intrinsic::nvvm_suld_3d_i8_trap:
case Intrinsic::nvvm_suld_3d_v2i8_trap:
case Intrinsic::nvvm_suld_3d_v4i8_trap:
case Intrinsic::nvvm_suld_1d_i8_zero:
case Intrinsic::nvvm_suld_1d_v2i8_zero:
case Intrinsic::nvvm_suld_1d_v4i8_zero:
case Intrinsic::nvvm_suld_1d_array_i8_zero:
case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
case Intrinsic::nvvm_suld_2d_i8_zero:
case Intrinsic::nvvm_suld_2d_v2i8_zero:
case Intrinsic::nvvm_suld_2d_v4i8_zero:
case Intrinsic::nvvm_suld_2d_array_i8_zero:
case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
case Intrinsic::nvvm_suld_3d_i8_zero:
case Intrinsic::nvvm_suld_3d_v2i8_zero:
case Intrinsic::nvvm_suld_3d_v4i8_zero:
Info.opc = getOpcForSurfaceInstr(Intrinsic);
Info.memVT = MVT::i8;
Info.ptrVal = nullptr;
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
case Intrinsic::nvvm_suld_1d_i16_clamp:
case Intrinsic::nvvm_suld_1d_v2i16_clamp:
case Intrinsic::nvvm_suld_1d_v4i16_clamp:
case Intrinsic::nvvm_suld_1d_array_i16_clamp:
case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
case Intrinsic::nvvm_suld_2d_i16_clamp:
case Intrinsic::nvvm_suld_2d_v2i16_clamp:
case Intrinsic::nvvm_suld_2d_v4i16_clamp:
case Intrinsic::nvvm_suld_2d_array_i16_clamp:
case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
case Intrinsic::nvvm_suld_3d_i16_clamp:
case Intrinsic::nvvm_suld_3d_v2i16_clamp:
case Intrinsic::nvvm_suld_3d_v4i16_clamp:
case Intrinsic::nvvm_suld_1d_i16_trap:
case Intrinsic::nvvm_suld_1d_v2i16_trap:
case Intrinsic::nvvm_suld_1d_v4i16_trap:
case Intrinsic::nvvm_suld_1d_array_i16_trap:
case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
case Intrinsic::nvvm_suld_2d_i16_trap:
case Intrinsic::nvvm_suld_2d_v2i16_trap:
case Intrinsic::nvvm_suld_2d_v4i16_trap:
case Intrinsic::nvvm_suld_2d_array_i16_trap:
case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
case Intrinsic::nvvm_suld_3d_i16_trap:
case Intrinsic::nvvm_suld_3d_v2i16_trap:
case Intrinsic::nvvm_suld_3d_v4i16_trap:
case Intrinsic::nvvm_suld_1d_i16_zero:
case Intrinsic::nvvm_suld_1d_v2i16_zero:
case Intrinsic::nvvm_suld_1d_v4i16_zero:
case Intrinsic::nvvm_suld_1d_array_i16_zero:
case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
case Intrinsic::nvvm_suld_2d_i16_zero:
case Intrinsic::nvvm_suld_2d_v2i16_zero:
case Intrinsic::nvvm_suld_2d_v4i16_zero:
case Intrinsic::nvvm_suld_2d_array_i16_zero:
case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
case Intrinsic::nvvm_suld_3d_i16_zero:
case Intrinsic::nvvm_suld_3d_v2i16_zero:
case Intrinsic::nvvm_suld_3d_v4i16_zero:
Info.opc = getOpcForSurfaceInstr(Intrinsic);
Info.memVT = MVT::i16;
Info.ptrVal = nullptr;
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
case Intrinsic::nvvm_suld_1d_i32_clamp:
case Intrinsic::nvvm_suld_1d_v2i32_clamp:
case Intrinsic::nvvm_suld_1d_v4i32_clamp:
case Intrinsic::nvvm_suld_1d_array_i32_clamp:
case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
case Intrinsic::nvvm_suld_2d_i32_clamp:
case Intrinsic::nvvm_suld_2d_v2i32_clamp:
case Intrinsic::nvvm_suld_2d_v4i32_clamp:
case Intrinsic::nvvm_suld_2d_array_i32_clamp:
case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
case Intrinsic::nvvm_suld_3d_i32_clamp:
case Intrinsic::nvvm_suld_3d_v2i32_clamp:
case Intrinsic::nvvm_suld_3d_v4i32_clamp:
case Intrinsic::nvvm_suld_1d_i32_trap:
case Intrinsic::nvvm_suld_1d_v2i32_trap:
case Intrinsic::nvvm_suld_1d_v4i32_trap:
case Intrinsic::nvvm_suld_1d_array_i32_trap:
case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
case Intrinsic::nvvm_suld_2d_i32_trap:
case Intrinsic::nvvm_suld_2d_v2i32_trap:
case Intrinsic::nvvm_suld_2d_v4i32_trap:
case Intrinsic::nvvm_suld_2d_array_i32_trap:
case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
case Intrinsic::nvvm_suld_3d_i32_trap:
case Intrinsic::nvvm_suld_3d_v2i32_trap:
case Intrinsic::nvvm_suld_3d_v4i32_trap:
case Intrinsic::nvvm_suld_1d_i32_zero:
case Intrinsic::nvvm_suld_1d_v2i32_zero:
case Intrinsic::nvvm_suld_1d_v4i32_zero:
case Intrinsic::nvvm_suld_1d_array_i32_zero:
case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
case Intrinsic::nvvm_suld_2d_i32_zero:
case Intrinsic::nvvm_suld_2d_v2i32_zero:
case Intrinsic::nvvm_suld_2d_v4i32_zero:
case Intrinsic::nvvm_suld_2d_array_i32_zero:
case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
case Intrinsic::nvvm_suld_3d_i32_zero:
case Intrinsic::nvvm_suld_3d_v2i32_zero:
case Intrinsic::nvvm_suld_3d_v4i32_zero:
Info.opc = getOpcForSurfaceInstr(Intrinsic);
Info.memVT = MVT::i32;
Info.ptrVal = nullptr;
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
case Intrinsic::nvvm_suld_1d_i64_clamp:
case Intrinsic::nvvm_suld_1d_v2i64_clamp:
case Intrinsic::nvvm_suld_1d_array_i64_clamp:
case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
case Intrinsic::nvvm_suld_2d_i64_clamp:
case Intrinsic::nvvm_suld_2d_v2i64_clamp:
case Intrinsic::nvvm_suld_2d_array_i64_clamp:
case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
case Intrinsic::nvvm_suld_3d_i64_clamp:
case Intrinsic::nvvm_suld_3d_v2i64_clamp:
case Intrinsic::nvvm_suld_1d_i64_trap:
case Intrinsic::nvvm_suld_1d_v2i64_trap:
case Intrinsic::nvvm_suld_1d_array_i64_trap:
case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
case Intrinsic::nvvm_suld_2d_i64_trap:
case Intrinsic::nvvm_suld_2d_v2i64_trap:
case Intrinsic::nvvm_suld_2d_array_i64_trap:
case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
case Intrinsic::nvvm_suld_3d_i64_trap:
case Intrinsic::nvvm_suld_3d_v2i64_trap:
case Intrinsic::nvvm_suld_1d_i64_zero:
case Intrinsic::nvvm_suld_1d_v2i64_zero:
case Intrinsic::nvvm_suld_1d_array_i64_zero:
case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
case Intrinsic::nvvm_suld_2d_i64_zero:
case Intrinsic::nvvm_suld_2d_v2i64_zero:
case Intrinsic::nvvm_suld_2d_array_i64_zero:
case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
case Intrinsic::nvvm_suld_3d_i64_zero:
case Intrinsic::nvvm_suld_3d_v2i64_zero:
Info.opc = getOpcForSurfaceInstr(Intrinsic);
Info.memVT = MVT::i64;
Info.ptrVal = nullptr;
Info.offset = 0;
Info.flags = MachineMemOperand::MOLoad;
Info.align = Align(16);
return true;
}
return false;
}