[llvm] r274664 - NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0
Justin Bogner via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 6 13:02:45 PDT 2016
Author: bogner
Date: Wed Jul 6 15:02:45 2016
New Revision: 274664
URL: http://llvm.org/viewvc/llvm-project?rev=274664&view=rev
Log:
NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0
Everywhere where cuda.syncthreads or __syncthreads is used, use the
properly namespaced nvvm.barrier0 instead.
Modified:
llvm/trunk/docs/LangRef.rst
llvm/trunk/docs/NVPTXUsage.rst
llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td
llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/trunk/test/CodeGen/NVPTX/MachineSink-convergent.ll
llvm/trunk/test/CodeGen/NVPTX/TailDuplication-convergent.ll
llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll
llvm/trunk/test/CodeGen/NVPTX/noduplicate-syncthreads.ll
llvm/trunk/test/Feature/intrinsic-noduplicate.ll
llvm/trunk/test/Transforms/FunctionAttrs/convergent.ll
Modified: llvm/trunk/docs/LangRef.rst
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/docs/LangRef.rst?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/docs/LangRef.rst (original)
+++ llvm/trunk/docs/LangRef.rst Wed Jul 6 15:02:45 2016
@@ -1318,7 +1318,7 @@ example:
The ``convergent`` attribute may appear on functions or call/invoke
instructions. When it appears on a function, it indicates that calls to
this function should not be made control-dependent on additional values.
- For example, the intrinsic ``llvm.cuda.syncthreads`` is ``convergent``, so
+ For example, the intrinsic ``llvm.nvvm.barrier0`` is ``convergent``, so
calls to this intrinsic cannot be made control-dependent on additional
values.
Modified: llvm/trunk/docs/NVPTXUsage.rst
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/docs/NVPTXUsage.rst?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/docs/NVPTXUsage.rst (original)
+++ llvm/trunk/docs/NVPTXUsage.rst Wed Jul 6 15:02:45 2016
@@ -566,7 +566,7 @@ Intrinsic
``i32 @llvm.nvvm.read.ptx.sreg.ctaid.{x,y,z}`` blockIdx.{x,y,z}
``i32 @llvm.nvvm.read.ptx.sreg.ntid.{x,y,z}`` blockDim.{x,y,z}
``i32 @llvm.nvvm.read.ptx.sreg.nctaid.{x,y,z}`` gridDim.{x,y,z}
-``void @llvm.cuda.syncthreads()`` __syncthreads()
+``void @llvm.nvvm.barrier0()`` __syncthreads()
================================================ ====================
Modified: llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td Wed Jul 6 15:02:45 2016
@@ -729,8 +729,6 @@ def llvm_anyi64ptr_ty : LLVMAnyPoint
[IntrArgMemOnly, NoCapture<0>]>;
// Bar.Sync
- def int_cuda_syncthreads : GCCBuiltin<"__syncthreads">,
- Intrinsic<[], [], [IntrConvergent]>;
def int_nvvm_barrier0 : GCCBuiltin<"__nvvm_bar0">,
Intrinsic<[], [], [IntrConvergent]>;
def int_nvvm_barrier0_popc : GCCBuiltin<"__nvvm_bar0_popc">,
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.cpp Wed Jul 6 15:02:45 2016
@@ -112,7 +112,7 @@ bool NVPTXInstrInfo::isStoreInstr(const
bool NVPTXInstrInfo::CanTailMerge(const MachineInstr *MI) const {
unsigned addrspace = 0;
- if (MI->getOpcode() == NVPTX::INT_CUDA_SYNCTHREADS)
+ if (MI->getOpcode() == NVPTX::INT_BARRIER0)
return false;
if (isLoadInstr(*MI, addrspace))
if (addrspace == NVPTX::PTXLdStInstCode::SHARED)
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td Wed Jul 6 15:02:45 2016
@@ -33,9 +33,6 @@ def immDouble1 : PatLeaf<(fpimm), [{
// Synchronization and shuffle functions
//-----------------------------------
let isConvergent = 1 in {
-def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
- "bar.sync \t0;",
- [(int_cuda_syncthreads)]>;
def INT_BARRIER0 : NVPTXInst<(outs), (ins),
"bar.sync \t0;",
[(int_nvvm_barrier0)]>;
Modified: llvm/trunk/test/CodeGen/NVPTX/MachineSink-convergent.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/MachineSink-convergent.ll?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/MachineSink-convergent.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/MachineSink-convergent.ll Wed Jul 6 15:02:45 2016
@@ -1,7 +1,7 @@
; RUN: llc < %s | FileCheck %s
target triple = "nvptx64-nvidia-cuda"
-declare void @llvm.cuda.syncthreads()
+declare void @llvm.nvvm.barrier0()
; 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
@@ -11,7 +11,7 @@ Start:
; CHECK: ld.u32
%ptr_val = load i32, i32* %ptr
; CHECK: bar.sync
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
br i1 %cond, label %L1, label %L2
L1:
%ptr_val2 = add i32 %ptr_val, 100
Modified: llvm/trunk/test/CodeGen/NVPTX/TailDuplication-convergent.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/TailDuplication-convergent.ll?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/TailDuplication-convergent.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/TailDuplication-convergent.ll Wed Jul 6 15:02:45 2016
@@ -2,7 +2,7 @@
target triple = "nvptx64-nvidia-cuda"
declare void @foo()
-declare void @llvm.cuda.syncthreads()
+declare void @llvm.nvvm.barrier0()
; syncthreads shouldn't be duplicated.
; CHECK: .func call_syncthreads
@@ -20,7 +20,7 @@ L2:
store i32 1, i32* %a
br label %L42
L42:
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
br label %Ret
}
Modified: llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/access-non-generic.ll Wed Jul 6 15:02:45 2016
@@ -34,7 +34,7 @@ define void @ld_st_shared_f32(i32 %i, fl
store float %v, float* addrspacecast (float addrspace(3)* @scalar to float*), align 4
; PTX: st.shared.f32 [scalar], %f{{[0-9]+}};
; use syncthreads to disable optimizations across components
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
; PTX: bar.sync 0;
; cast; load
@@ -45,7 +45,7 @@ define void @ld_st_shared_f32(i32 %i, fl
; cast; store
store float %v, float* %2, align 4
; PTX: st.shared.f32 [scalar], %f{{[0-9]+}};
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
; PTX: bar.sync 0;
; load gep cast
@@ -55,7 +55,7 @@ define void @ld_st_shared_f32(i32 %i, fl
; store gep cast
store float %v, float* getelementptr inbounds ([10 x float], [10 x float]* addrspacecast ([10 x float] addrspace(3)* @array to [10 x float]*), i32 0, i32 5), align 4
; PTX: st.shared.f32 [array+20], %f{{[0-9]+}};
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
; PTX: bar.sync 0;
; gep cast; load
@@ -66,7 +66,7 @@ define void @ld_st_shared_f32(i32 %i, fl
; gep cast; store
store float %v, float* %5, align 4
; PTX: st.shared.f32 [array+20], %f{{[0-9]+}};
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
; PTX: bar.sync 0;
; cast; gep; load
@@ -78,7 +78,7 @@ define void @ld_st_shared_f32(i32 %i, fl
; cast; gep; store
store float %v, float* %8, align 4
; PTX: st.shared.f32 [%{{(r|rl|rd)[0-9]+}}], %f{{[0-9]+}};
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
; PTX: bar.sync 0;
ret void
@@ -181,7 +181,7 @@ exit:
ret void
}
-declare void @llvm.cuda.syncthreads() #3
+declare void @llvm.nvvm.barrier0() #3
declare void @use(float)
Modified: llvm/trunk/test/CodeGen/NVPTX/noduplicate-syncthreads.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/noduplicate-syncthreads.ll?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/noduplicate-syncthreads.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/noduplicate-syncthreads.ll Wed Jul 6 15:02:45 2016
@@ -3,8 +3,8 @@
; Make sure the call to syncthreads is not duplicate here by the LLVM
; optimizations, because it has the noduplicate attribute set.
-; CHECK: call void @llvm.cuda.syncthreads
-; CHECK-NOT: call void @llvm.cuda.syncthreads
+; CHECK: call void @llvm.nvvm.barrier0
+; CHECK-NOT: call void @llvm.nvvm.barrier0
; Function Attrs: nounwind
define void @foo(float* %output) #1 {
@@ -37,7 +37,7 @@ if.else:
br label %if.end
if.end: ; preds = %if.else, %if.then
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
%6 = load float*, float** %output.addr, align 8
%arrayidx6 = getelementptr inbounds float, float* %6, i64 0
%7 = load float, float* %arrayidx6, align 4
@@ -68,7 +68,7 @@ if.end17:
}
; Function Attrs: noduplicate nounwind
-declare void @llvm.cuda.syncthreads() #2
+declare void @llvm.nvvm.barrier0() #2
!0 = !{void (float*)* @foo, !"kernel", i32 1}
!1 = !{null, !"align", i32 8}
Modified: llvm/trunk/test/Feature/intrinsic-noduplicate.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Feature/intrinsic-noduplicate.ll?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/test/Feature/intrinsic-noduplicate.ll (original)
+++ llvm/trunk/test/Feature/intrinsic-noduplicate.ll Wed Jul 6 15:02:45 2016
@@ -1,9 +1,9 @@
; RUN: llvm-as < %s | llvm-dis | FileCheck %s
; Make sure LLVM knows about the convergent attribute on the
-; llvm.cuda.syncthreads intrinsic.
+; llvm.nvvm.barrier0 intrinsic.
-declare void @llvm.cuda.syncthreads()
+declare void @llvm.nvvm.barrier0()
-; CHECK: declare void @llvm.cuda.syncthreads() #[[ATTRNUM:[0-9]+]]
+; CHECK: declare void @llvm.nvvm.barrier0() #[[ATTRNUM:[0-9]+]]
; CHECK: attributes #[[ATTRNUM]] = { convergent nounwind }
Modified: llvm/trunk/test/Transforms/FunctionAttrs/convergent.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/Transforms/FunctionAttrs/convergent.ll?rev=274664&r1=274663&r2=274664&view=diff
==============================================================================
--- llvm/trunk/test/Transforms/FunctionAttrs/convergent.ll (original)
+++ llvm/trunk/test/Transforms/FunctionAttrs/convergent.ll Wed Jul 6 15:02:45 2016
@@ -59,15 +59,15 @@ define i32 @indirect_non_convergent_call
; CHECK: Function Attrs
; CHECK-SAME: convergent
-; CHECK-NEXT: declare void @llvm.cuda.syncthreads()
-declare void @llvm.cuda.syncthreads() convergent
+; CHECK-NEXT: declare void @llvm.nvvm.barrier0()
+declare void @llvm.nvvm.barrier0() convergent
; CHECK: Function Attrs
; CHECK-SAME: convergent
; CHECK-NEXT: define i32 @intrinsic()
define i32 @intrinsic() convergent {
; Implicitly convergent, because the intrinsic is convergent.
- call void @llvm.cuda.syncthreads()
+ call void @llvm.nvvm.barrier0()
ret i32 0
}
More information about the llvm-commits
mailing list