This is an archive of the discontinued LLVM Phabricator instance.

Differential D123039

[NVPTX] shfl.sync is introduced in PTX 6.0
ClosedPublic

Authored by asavonic on Apr 4 2022, 8:35 AM.

Details

Reviewers
tra
jholewinski
Commits
rG230f32696497: [NVPTX] shfl.sync is introduced in PTX 6.0
Summary

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.

Diff Detail

Event Timeline

asavonic created this revision.Apr 4 2022, 8:35 AM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2022, 8:35 AM
Herald added a subscriber: hiraditya. · View Herald Transcript
asavonic requested review of this revision.Apr 4 2022, 8:35 AM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2022, 8:35 AM
Herald added subscribers: llvm-commits, jdoerfert. · View Herald Transcript
asavonic mentioned this in D121727: [NVPTX] Integrate ptxas to LIT tests.Apr 4 2022, 8:36 AM
Harbormaster completed remote builds in B157743: Diff 420172.Apr 4 2022, 9:10 AM
tra accepted this revision.Apr 4 2022, 11:03 AM
This revision is now accepted and ready to land.Apr 4 2022, 11:03 AM
asavonic added a parent revision: D121727: [NVPTX] Integrate ptxas to LIT tests.Apr 11 2022, 8:52 AM
This revision was landed with ongoing or failed builds.Apr 14 2022, 7:08 AM
Closed by commit rG230f32696497: [NVPTX] shfl.sync is introduced in PTX 6.0 (authored by asavonic). · Explain Why
This revision was automatically updated to reflect the committed changes.
asavonic added a commit: rG230f32696497: [NVPTX] shfl.sync is introduced in PTX 6.0.
Herald added a subscriber: gchakrabarti. · View Herald TranscriptApr 14 2022, 7:08 AM

Revision Contents

PathSize
llvm/
lib/
Target/
NVPTX/
2 lines
test/
CodeGen/
NVPTX/
2 lines
2 lines

Diff 420172

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 176 Lines • ▼ Show 20 Linesforeach sync = [false, true] in {
foreach mode = ["up", "down", "bfly", "idx"] in { foreach mode = ["up", "down", "bfly", "idx"] in {
foreach regclass = ["i32", "f32"] in { foreach regclass = ["i32", "f32"] in {
foreach return_pred = [false, true] in { foreach return_pred = [false, true] in {
foreach offset_imm = [false, true] in { foreach offset_imm = [false, true] in {
foreach mask_imm = [false, true] in { foreach mask_imm = [false, true] in {
foreach threadmask_imm = THREADMASK_INFO<sync>.ret in { foreach threadmask_imm = THREADMASK_INFO<sync>.ret in {
def : SHFL_INSTR<sync, mode, regclass, return_pred, def : SHFL_INSTR<sync, mode, regclass, return_pred,
offset_imm, mask_imm, threadmask_imm>, offset_imm, mask_imm, threadmask_imm>,
Requires<!if(sync, [hasSM30], [hasSM30, hasSHFL])>; Requires<!if(sync, [hasSM30, hasPTX60], [hasSM30, hasSHFL])>;
} }
} }
} }
} }
} }
} }
} }
▲ Show 20 Lines • Show All 6,501 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/shfl-sync-p.ll

; 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 {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) declare {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32, float, i32, i32)
declare {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32) declare {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32)
declare {float, i1} @llvm.nvvm.shfl.sync.up.f32p(i32, float, i32, i32) declare {float, i1} @llvm.nvvm.shfl.sync.up.f32p(i32, float, i32, i32)
declare {i32, i1} @llvm.nvvm.shfl.sync.bfly.i32p(i32, i32, i32, i32) declare {i32, i1} @llvm.nvvm.shfl.sync.bfly.i32p(i32, i32, i32, i32)
declare {float, i1} @llvm.nvvm.shfl.sync.bfly.f32p(i32, float, i32, i32) declare {float, i1} @llvm.nvvm.shfl.sync.bfly.f32p(i32, float, i32, i32)
declare {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32) declare {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32)
▲ Show 20 Lines • Show All 171 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/shfl-sync.ll

; 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 i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32) declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32) declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32) declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32) declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32) declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32) declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
▲ Show 20 Lines • Show All 85 Lines • Show Last 20 Lines