[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