This is an archive of the discontinued LLVM Phabricator instance.

Differential D17319

[NVPTX] Test that MachineSink won't sink across llvm.cuda.syncthreads.
ClosedPublic

Authored by jlebar on Feb 16 2016, 5:03 PM.

Details

Reviewers
jingyue
Commits
rGb5c7b1c00fc3: [NVPTX] Test that MachineSink won't sink across llvm.cuda.syncthreads.
rL261114: [NVPTX] Test that MachineSink won't sink across llvm.cuda.syncthreads.
Summary

The syncthreads MI is modeled as mayread/maywrite -- convergence doesn't
even come into play here. Nonetheless this property is highly implicit
in the tablegen files, so a test seems appropriate.

Diff Detail

Event Timeline

jlebar updated this revision to Diff 48134.Feb 16 2016, 5:03 PM
jlebar retitled this revision from to [NVPTX] Test that MachineSink won't sink across llvm.cuda.syncthreads..
jlebar updated this object.
jlebar added a reviewer: jingyue.
jlebar added a subscriber: llvm-commits.
Herald added a subscriber: jholewinski. · View Herald TranscriptFeb 16 2016, 5:03 PM
jingyue accepted this revision.Feb 16 2016, 8:53 PM
jingyue edited edge metadata.
This revision is now accepted and ready to land.Feb 16 2016, 8:53 PM
Closed by commit rL261114: [NVPTX] Test that MachineSink won't sink across llvm.cuda.syncthreads. (authored by jlebar). · Explain WhyFeb 17 2016, 9:51 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
test/
CodeGen/
NVPTX/
23 lines

Diff 48134

test/CodeGen/NVPTX/MachineSink-convergent.ll

  • This file was added.
; RUN: llc < %s | FileCheck %s
target triple = "nvptx64-nvidia-cuda"
declare void @llvm.cuda.syncthreads()
; Load a value, then syncthreads. Branch, and use the loaded value only on one
; side of the branch. The load shouldn't be sunk beneath the call, because
; syncthreads is modeled as maystore.
define i32 @f(i32 %x, i32* %ptr, i1 %cond) {
Start:
; CHECK: ld.u32
%ptr_val = load i32, i32* %ptr
; CHECK: bar.sync
call void @llvm.cuda.syncthreads()
br i1 %cond, label %L1, label %L2
L1:
%ptr_val2 = add i32 %ptr_val, 100
br label %L2
L2:
%v4 = phi i32 [ %x, %Start ], [ %ptr_val2, %L1 ]
%v5 = add i32 %v4, 1000
ret i32 %v5
}