[llvm] [LLVM][NVPTX] Add codegen support for tcgen05.{ld, st} instructions (PR #126740)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 11 07:09:04 PST 2025
github-actions[bot] wrote:
<!--LLVM CODE FORMAT COMMENT: {clang-format}-->
:warning: C/C++ code formatter, clang-format found issues in your code. :warning:
<details>
<summary>
You can test this locally with the following command:
</summary>
``````````bash
git-clang-format --diff d89c23b487dfb96718f9f72f2841a83906cc3d0b 56e2a96146f4b6639e227561830c681050ca5ce7 --extensions h,cpp -- llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
``````````
</details>
<details>
<summary>
View the diff from clang-format here.
</summary>
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 4a51316c2b..1cf5a9b7b9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -3682,7 +3682,7 @@ bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32:
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64:
case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128: {
- SelectTcgen05St(N,/* hasOffset */ true);
+ SelectTcgen05St(N, /* hasOffset */ true);
return true;
}
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 5057c849ba..bf8127990b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -4340,8 +4340,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x64b_x1:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x1:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x1:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x1: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v1i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4354,8 +4353,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x64b_x2:
case Intrinsic::nvvm_tcgen05_ld_16x128b_x1:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x2:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x2:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x2: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v2i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4369,8 +4367,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x128b_x2:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x4:
case Intrinsic::nvvm_tcgen05_ld_16x256b_x1:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x4:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x4: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v4i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4384,8 +4381,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x128b_x4:
case Intrinsic::nvvm_tcgen05_ld_16x256b_x2:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x8:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x8:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x8: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v8i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4399,8 +4395,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x128b_x8:
case Intrinsic::nvvm_tcgen05_ld_16x256b_x4:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x16:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x16:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x16: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v16i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4414,8 +4409,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x128b_x16:
case Intrinsic::nvvm_tcgen05_ld_16x256b_x8:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x32:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x32:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x32: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v32i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4429,8 +4423,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x128b_x32:
case Intrinsic::nvvm_tcgen05_ld_16x256b_x16:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x64:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x64:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x64: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v64i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4444,8 +4437,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_ld_16x128b_x64:
case Intrinsic::nvvm_tcgen05_ld_16x256b_x32:
case Intrinsic::nvvm_tcgen05_ld_32x32b_x128:
- case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x128:
- {
+ case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x128: {
Info.opc = ISD::INTRINSIC_W_CHAIN;
Info.memVT = MVT::v128i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4457,8 +4449,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x64b_x1:
case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4471,8 +4462,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
case Intrinsic::nvvm_tcgen05_st_16x128b_x1:
case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v2i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4486,8 +4476,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x128b_x2:
case Intrinsic::nvvm_tcgen05_st_16x256b_x1:
case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v4i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4501,8 +4490,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x128b_x4:
case Intrinsic::nvvm_tcgen05_st_16x256b_x2:
case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v8i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4516,8 +4504,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x128b_x8:
case Intrinsic::nvvm_tcgen05_st_16x256b_x4:
case Intrinsic::nvvm_tcgen05_st_32x32b_x16:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v16i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4531,8 +4518,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x128b_x16:
case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
case Intrinsic::nvvm_tcgen05_st_32x32b_x32:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v32i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4546,8 +4532,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x128b_x32:
case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v64i32;
Info.ptrVal = I.getArgOperand(0);
@@ -4561,8 +4546,7 @@ bool NVPTXTargetLowering::getTgtMemIntrinsic(
case Intrinsic::nvvm_tcgen05_st_16x128b_x64:
case Intrinsic::nvvm_tcgen05_st_16x256b_x32:
case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
- case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128:
- {
+ case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128: {
Info.opc = ISD::INTRINSIC_VOID;
Info.memVT = MVT::v128i32;
Info.ptrVal = I.getArgOperand(0);
``````````
</details>
https://github.com/llvm/llvm-project/pull/126740
More information about the llvm-commits
mailing list