[PATCH] D123039: [NVPTX] shfl.sync is introduced in PTX 6.0

[Andrew Savonichev via Phabricator via llvm-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG230f32696497: [NVPTX] shfl.sync is introduced in PTX 6.0 (authored by asavonic).
Herald added a subscriber: gchakrabarti.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D123039/new/

https://reviews.llvm.org/D123039

Files:
  llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
  llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
  llvm/test/CodeGen/NVPTX/shfl-sync.ll


Index: llvm/test/CodeGen/NVPTX/shfl-sync.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/shfl-sync.ll
+++ llvm/test/CodeGen/NVPTX/shfl-sync.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
 
 declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
 declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
Index: llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
+++ llvm/test/CodeGen/NVPTX/shfl-sync-p.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
 
 declare {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32, i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32, float, i32, i32)
Index: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -182,7 +182,7 @@
             foreach threadmask_imm = THREADMASK_INFO<sync>.ret in {
               def : SHFL_INSTR<sync, mode, regclass, return_pred,
                                offset_imm, mask_imm, threadmask_imm>,
-                    Requires<!if(sync, [hasSM30], [hasSM30, hasSHFL])>;
+                    Requires<!if(sync, [hasSM30, hasPTX60], [hasSM30, hasSHFL])>;
             }
           }
         }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D123039.422855.patch
Type: text/x-patch
Size: 1596 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220414/7cb588b5/attachment.bin>