diff options
| author | Justin Bogner <mail@justinbogner.com> | 2016-07-06 20:02:45 +0000 |
|---|---|---|
| committer | Justin Bogner <mail@justinbogner.com> | 2016-07-06 20:02:45 +0000 |
| commit | a463537a3644d4013b23be20e5446af609342f01 (patch) | |
| tree | fbd009adc55df065612b62a6727a50cfae640225 /llvm/lib/Target | |
| parent | 2f8de9fb4fea7bd29ffd2ac85e9a0f20ea1410ca (diff) | |
| download | bcm5719-llvm-a463537a3644d4013b23be20e5446af609342f01.tar.gz bcm5719-llvm-a463537a3644d4013b23be20e5446af609342f01.zip | |
NVPTX: Replace uses of cuda.syncthreads with nvvm.barrier0
Everywhere where cuda.syncthreads or __syncthreads is used, use the
properly namespaced nvvm.barrier0 instead.
llvm-svn: 274664
Diffstat (limited to 'llvm/lib/Target')
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp | 2 | ||||
| -rw-r--r-- | llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 3 |
2 files changed, 1 insertions, 4 deletions
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp index 3cebfce0f93..cd1d355937f 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp @@ -112,7 +112,7 @@ bool NVPTXInstrInfo::isStoreInstr(const MachineInstr &MI, 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) diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td index 6f7df39c771..6e40421323e 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -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)]>; |
