summaryrefslogtreecommitdiffstats
path: root/llvm
diff options
context:
space:
mode:
Diffstat (limited to 'llvm')
-rw-r--r--llvm/lib/Target/NVPTX/CMakeLists.txt1
-rw-r--r--llvm/lib/Target/NVPTX/NVPTX.h1
-rw-r--r--llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp3
-rw-r--r--llvm/lib/Target/NVPTX/NVVMIntrRange.cpp154
-rw-r--r--llvm/test/CodeGen/NVPTX/intrinsic-old.ll67
5 files changed, 213 insertions, 13 deletions
diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt
index ed43296b870..b67c4050086 100644
--- a/llvm/lib/Target/NVPTX/CMakeLists.txt
+++ b/llvm/lib/Target/NVPTX/CMakeLists.txt
@@ -32,6 +32,7 @@ set(NVPTXCodeGen_sources
NVPTXTargetMachine.cpp
NVPTXTargetTransformInfo.cpp
NVPTXUtilities.cpp
+ NVVMIntrRange.cpp
NVVMReflect.cpp
)
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 15a4205f6aa..e91385ac13f 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -47,6 +47,7 @@ ModulePass *createNVPTXAssignValidGlobalNamesPass();
ModulePass *createGenericToNVVMPass();
FunctionPass *createNVPTXFavorNonGenericAddrSpacesPass();
FunctionPass *createNVPTXInferAddressSpacesPass();
+FunctionPass *createNVVMIntrRangePass(unsigned int SmVersion);
FunctionPass *createNVVMReflectPass();
FunctionPass *createNVVMReflectPass(const StringMap<int> &Mapping);
MachineFunctionPass *createNVPTXPrologEpilogPass();
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index 583baad65c9..b9f5919964c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -55,6 +55,7 @@ static cl::opt<bool> UseInferAddressSpaces(
"NVPTXFavorNonGenericAddrSpaces"));
namespace llvm {
+void initializeNVVMIntrRangePass(PassRegistry&);
void initializeNVVMReflectPass(PassRegistry&);
void initializeGenericToNVVMPass(PassRegistry&);
void initializeNVPTXAllocaHoistingPass(PassRegistry &);
@@ -75,6 +76,7 @@ extern "C" void LLVMInitializeNVPTXTarget() {
// but it's very NVPTX-specific.
PassRegistry &PR = *PassRegistry::getPassRegistry();
initializeNVVMReflectPass(PR);
+ initializeNVVMIntrRangePass(PR);
initializeGenericToNVVMPass(PR);
initializeNVPTXAllocaHoistingPass(PR);
initializeNVPTXAssignValidGlobalNamesPass(PR);
@@ -176,6 +178,7 @@ TargetPassConfig *NVPTXTargetMachine::createPassConfig(PassManagerBase &PM) {
void NVPTXTargetMachine::addEarlyAsPossiblePasses(PassManagerBase &PM) {
PM.add(createNVVMReflectPass());
+ PM.add(createNVVMIntrRangePass(Subtarget.getSmVersion()));
}
TargetIRAnalysis NVPTXTargetMachine::getTargetIRAnalysis() {
diff --git a/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp b/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp
new file mode 100644
index 00000000000..09f328d4c53
--- /dev/null
+++ b/llvm/lib/Target/NVPTX/NVVMIntrRange.cpp
@@ -0,0 +1,154 @@
+//===- NVVMIntrRange.cpp - Set !range metadata for NVVM intrinsics --------===//
+//
+// The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This pass adds appropriate !range metadata for calls to NVVM
+// intrinsics that return a limited range of values.
+//
+//===----------------------------------------------------------------------===//
+
+#include "NVPTX.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/InstIterator.h"
+#include "llvm/IR/Intrinsics.h"
+#include "llvm/IR/Instructions.h"
+
+using namespace llvm;
+
+#define DEBUG_TYPE "nvvm-intr-range"
+
+namespace llvm { void initializeNVVMIntrRangePass(PassRegistry &); }
+
+// Add !range metadata based on limits of given SM variant.
+static cl::opt<unsigned> NVVMIntrRangeSM("nvvm-intr-range-sm", cl::init(20),
+ cl::Hidden, cl::desc("SM variant"));
+
+namespace {
+class NVVMIntrRange : public FunctionPass {
+ private:
+ struct {
+ unsigned x, y, z;
+ } MaxBlockSize, MaxGridSize;
+
+ public:
+ static char ID;
+ NVVMIntrRange() : NVVMIntrRange(NVVMIntrRangeSM) {}
+ NVVMIntrRange(unsigned int SmVersion)
+ : FunctionPass(ID), MaxBlockSize{1024, 1024, 64},
+ MaxGridSize{SmVersion >= 30 ? 0x7fffffffu : 0xffffu, 0xffff, 0xffff} {
+ initializeNVVMIntrRangePass(*PassRegistry::getPassRegistry());
+ }
+
+ bool runOnFunction(Function &) override;
+};
+}
+
+FunctionPass *llvm::createNVVMIntrRangePass(unsigned int SmVersion) {
+ return new NVVMIntrRange(SmVersion);
+}
+
+char NVVMIntrRange::ID = 0;
+INITIALIZE_PASS(NVVMIntrRange, "nvvm-intr-range",
+ "Add !range metadata to NVVM intrinsics.", false, false)
+
+// Adds the passed-in [Low,High) range information as metadata to the
+// passed-in call instruction.
+static bool addRangeMetadata(uint64_t Low, uint64_t High, CallInst *C) {
+ LLVMContext &Context = C->getParent()->getContext();
+ IntegerType *Int32Ty = Type::getInt32Ty(Context);
+ Metadata *LowAndHigh[] = {
+ ConstantAsMetadata::get(ConstantInt::get(Int32Ty, Low)),
+ ConstantAsMetadata::get(ConstantInt::get(Int32Ty, High))};
+ C->setMetadata(LLVMContext::MD_range, MDNode::get(Context, LowAndHigh));
+ return true;
+}
+
+bool NVVMIntrRange::runOnFunction(Function &F) {
+ // Go through the calls in this function.
+ bool Changed = false;
+ for (Instruction &I : instructions(F)) {
+ CallInst *Call = dyn_cast<CallInst>(&I);
+ if (!Call)
+ continue;
+
+ if (Function *Callee = Call->getCalledFunction()) {
+ switch (Callee->getIntrinsicID()) {
+ // Index within block
+ case Intrinsic::ptx_read_tid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_x:
+ Changed |= addRangeMetadata(0, MaxBlockSize.x, Call);
+ break;
+ case Intrinsic::ptx_read_tid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_y:
+ Changed |= addRangeMetadata(0, MaxBlockSize.y, Call);
+ break;
+ case Intrinsic::ptx_read_tid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_z:
+ Changed |= addRangeMetadata(0, MaxBlockSize.z, Call);
+ break;
+
+ // Block size
+ case Intrinsic::ptx_read_ntid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_x:
+ Changed |= addRangeMetadata(1, MaxBlockSize.x+1, Call);
+ break;
+ case Intrinsic::ptx_read_ntid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_y:
+ Changed |= addRangeMetadata(1, MaxBlockSize.y+1, Call);
+ break;
+ case Intrinsic::ptx_read_ntid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_ntid_z:
+ Changed |= addRangeMetadata(1, MaxBlockSize.z+1, Call);
+ break;
+
+ // Index within grid
+ case Intrinsic::ptx_read_ctaid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_x:
+ Changed |= addRangeMetadata(0, MaxGridSize.x, Call);
+ break;
+ case Intrinsic::ptx_read_ctaid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_y:
+ Changed |= addRangeMetadata(0, MaxGridSize.y, Call);
+ break;
+ case Intrinsic::ptx_read_ctaid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_ctaid_z:
+ Changed |= addRangeMetadata(0, MaxGridSize.z, Call);
+ break;
+
+ // Grid size
+ case Intrinsic::ptx_read_nctaid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_x:
+ Changed |= addRangeMetadata(1, MaxGridSize.x+1, Call);
+ break;
+ case Intrinsic::ptx_read_nctaid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_y:
+ Changed |= addRangeMetadata(1, MaxGridSize.y+1, Call);
+ break;
+ case Intrinsic::ptx_read_nctaid_z:
+ case Intrinsic::nvvm_read_ptx_sreg_nctaid_z:
+ Changed |= addRangeMetadata(1, MaxGridSize.z+1, Call);
+ break;
+
+ // warp size is constant 32.
+ case Intrinsic::nvvm_read_ptx_sreg_warpsize:
+ Changed |= addRangeMetadata(32, 32+1, Call);
+ break;
+
+ // Lane ID is [0..warpsize)
+ case Intrinsic::ptx_read_laneid:
+ Changed |= addRangeMetadata(0, 32, Call);
+ break;
+
+ default:
+ break;
+ }
+ }
+ }
+
+ return Changed;
+}
diff --git a/llvm/test/CodeGen/NVPTX/intrinsic-old.ll b/llvm/test/CodeGen/NVPTX/intrinsic-old.ll
index 3c51776c0ec..5c73f44d075 100644
--- a/llvm/test/CodeGen/NVPTX/intrinsic-old.ll
+++ b/llvm/test/CodeGen/NVPTX/intrinsic-old.ll
@@ -1,8 +1,14 @@
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
+; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda -nvvm-intr-range \
+; RUN: | FileCheck --check-prefix=RANGE --check-prefix=RANGE_20 %s
+; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda \
+; RUN: -nvvm-intr-range -nvvm-intr-range-sm=30 \
+; RUN: | FileCheck --check-prefix=RANGE --check-prefix=RANGE_30 %s
define ptx_device i32 @test_tid_x() {
; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x;
+; RANGE: call i32 @llvm.ptx.read.tid.x(), !range ![[BLK_IDX_XY:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.tid.x()
ret i32 %x
@@ -10,6 +16,7 @@ define ptx_device i32 @test_tid_x() {
define ptx_device i32 @test_tid_y() {
; CHECK: mov.u32 %r{{[0-9]+}}, %tid.y;
+; RANGE: call i32 @llvm.ptx.read.tid.y(), !range ![[BLK_IDX_XY]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.tid.y()
ret i32 %x
@@ -17,6 +24,7 @@ define ptx_device i32 @test_tid_y() {
define ptx_device i32 @test_tid_z() {
; CHECK: mov.u32 %r{{[0-9]+}}, %tid.z;
+; RANGE: call i32 @llvm.ptx.read.tid.z(), !range ![[BLK_IDX_Z:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.tid.z()
ret i32 %x
@@ -31,6 +39,7 @@ define ptx_device i32 @test_tid_w() {
define ptx_device i32 @test_ntid_x() {
; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.x;
+; RANGE: call i32 @llvm.ptx.read.ntid.x(), !range ![[BLK_SIZE_XY:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.ntid.x()
ret i32 %x
@@ -38,6 +47,7 @@ define ptx_device i32 @test_ntid_x() {
define ptx_device i32 @test_ntid_y() {
; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.y;
+; RANGE: call i32 @llvm.ptx.read.ntid.y(), !range ![[BLK_SIZE_XY]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.ntid.y()
ret i32 %x
@@ -45,6 +55,7 @@ define ptx_device i32 @test_ntid_y() {
define ptx_device i32 @test_ntid_z() {
; CHECK: mov.u32 %r{{[0-9]+}}, %ntid.z;
+; RANGE: call i32 @llvm.ptx.read.ntid.z(), !range ![[BLK_SIZE_Z:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.ntid.z()
ret i32 %x
@@ -59,11 +70,20 @@ define ptx_device i32 @test_ntid_w() {
define ptx_device i32 @test_laneid() {
; CHECK: mov.u32 %r{{[0-9]+}}, %laneid;
+; RANGE: call i32 @llvm.ptx.read.laneid(), !range ![[LANEID:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.laneid()
ret i32 %x
}
+define ptx_device i32 @test_warpsize() {
+; CHECK: mov.u32 %r{{[0-9]+}}, WARP_SZ;
+; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.warpsize(), !range ![[WARPSIZE:[0-9]+]]
+; CHECK: ret;
+ %x = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
+ ret i32 %x
+}
+
define ptx_device i32 @test_warpid() {
; CHECK: mov.u32 %r{{[0-9]+}}, %warpid;
; CHECK: ret;
@@ -78,15 +98,9 @@ define ptx_device i32 @test_nwarpid() {
ret i32 %x
}
-define ptx_device i32 @test_ctaid_x() {
-; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x;
-; CHECK: ret;
- %x = call i32 @llvm.ptx.read.ctaid.x()
- ret i32 %x
-}
-
define ptx_device i32 @test_ctaid_y() {
; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.y;
+; RANGE: call i32 @llvm.ptx.read.ctaid.y(), !range ![[GRID_IDX_YZ:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.ctaid.y()
ret i32 %x
@@ -94,27 +108,31 @@ define ptx_device i32 @test_ctaid_y() {
define ptx_device i32 @test_ctaid_z() {
; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.z;
+; RANGE: call i32 @llvm.ptx.read.ctaid.z(), !range ![[GRID_IDX_YZ]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.ctaid.z()
ret i32 %x
}
-define ptx_device i32 @test_ctaid_w() {
-; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w;
+define ptx_device i32 @test_ctaid_x() {
+; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.x;
+; RANGE_30: call i32 @llvm.ptx.read.ctaid.x(), !range ![[GRID_IDX_X:[0-9]+]]
+; RANGE_20: call i32 @llvm.ptx.read.ctaid.x(), !range ![[GRID_IDX_YZ]]
; CHECK: ret;
- %x = call i32 @llvm.ptx.read.ctaid.w()
+ %x = call i32 @llvm.ptx.read.ctaid.x()
ret i32 %x
}
-define ptx_device i32 @test_nctaid_x() {
-; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;
+define ptx_device i32 @test_ctaid_w() {
+; CHECK: mov.u32 %r{{[0-9]+}}, %ctaid.w;
; CHECK: ret;
- %x = call i32 @llvm.ptx.read.nctaid.x()
+ %x = call i32 @llvm.ptx.read.ctaid.w()
ret i32 %x
}
define ptx_device i32 @test_nctaid_y() {
; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.y;
+; RANGE: call i32 @llvm.ptx.read.nctaid.y(), !range ![[GRID_SIZE_YZ:[0-9]+]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.nctaid.y()
ret i32 %x
@@ -122,11 +140,22 @@ define ptx_device i32 @test_nctaid_y() {
define ptx_device i32 @test_nctaid_z() {
; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.z;
+; RANGE: call i32 @llvm.ptx.read.nctaid.z(), !range ![[GRID_SIZE_YZ]]
; CHECK: ret;
%x = call i32 @llvm.ptx.read.nctaid.z()
ret i32 %x
}
+define ptx_device i32 @test_nctaid_x() {
+; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.x;
+; RANGE_30: call i32 @llvm.ptx.read.nctaid.x(), !range ![[GRID_SIZE_X:[0-9]+]]
+; RANGE_20: call i32 @llvm.ptx.read.nctaid.x(), !range ![[GRID_SIZE_YZ]]
+; CHECK: ret;
+ %x = call i32 @llvm.ptx.read.nctaid.x()
+ ret i32 %x
+}
+
+
define ptx_device i32 @test_nctaid_w() {
; CHECK: mov.u32 %r{{[0-9]+}}, %nctaid.w;
; CHECK: ret;
@@ -248,6 +277,7 @@ declare i32 @llvm.ptx.read.ntid.y()
declare i32 @llvm.ptx.read.ntid.z()
declare i32 @llvm.ptx.read.ntid.w()
+declare i32 @llvm.nvvm.read.ptx.sreg.warpsize()
declare i32 @llvm.ptx.read.laneid()
declare i32 @llvm.ptx.read.warpid()
declare i32 @llvm.ptx.read.nwarpid()
@@ -280,3 +310,14 @@ declare i32 @llvm.ptx.read.pm2()
declare i32 @llvm.ptx.read.pm3()
declare void @llvm.ptx.bar.sync(i32 %i)
+
+; RANGE-DAG: ![[BLK_IDX_XY]] = !{i32 0, i32 1024}
+; RANGE-DAG: ![[BLK_IDX_Z]] = !{i32 0, i32 64}
+; RANGE-DAG: ![[BLK_SIZE_XY]] = !{i32 1, i32 1025}
+; RANGE-DAG: ![[BLK_SIZE_Z]] = !{i32 1, i32 65}
+; RANGE-DAG: ![[LANEID]] = !{i32 0, i32 32}
+; RANGE-DAG: ![[WARPSIZE]] = !{i32 32, i32 33}
+; RANGE_30-DAG: ![[GRID_IDX_X]] = !{i32 0, i32 2147483647}
+; RANGE-DAG: ![[GRID_IDX_YZ]] = !{i32 0, i32 65535}
+; RANGE_30-DAG: ![[GRID_SIZE_X]] = !{i32 1, i32 -2147483648}
+; RANGE-DAG: ![[GRID_SIZE_YZ]] = !{i32 1, i32 65536}
OpenPOWER on IntegriCloud