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

[Andrew Savonichev via Phabricator via llvm-commits]

asavonic created this revision.
asavonic added reviewers: tra, jholewinski.
Herald added a subscriber: hiraditya.
Herald added a project: All.
asavonic requested review of this revision.
Herald added subscribers: llvm-commits, jdoerfert.
Herald added a project: LLVM.

PTX ISA spec, s9.7.8.6. Data Movement and Conversion Instructions: shfl.sync

PTX ISA Notes
Introduced in PTX ISA version 6.0.

Target ISA Notes
Requires sm_30 or higher.


Repository:
  rG LLVM Github Monorepo

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.420172.patch
Type: text/x-patch
Size: 1596 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220404/a94fb0e6/attachment.bin>