Differential D38978 Diff 129093 lib/Target/NVPTX/NVPTXFunctionDataSharing.cpp

Changeset View

Standalone View

lib/Target/NVPTX/NVPTXFunctionDataSharing.cpp

This file was added.

				//===-- FunctionDataSharing.cpp - Mark pointers as shared -----------------===//
				//
				// The LLVM Compiler Infrastructure
				//
				// This file is distributed under the University of Illinois Open Source
				// License. See LICENSE.TXT for details.
				//
				//===----------------------------------------------------------------------===//
				//
				// For all alloca instructions, add a pair of cast to shared address for
				traUnsubmitted Done Reply Inline Actions Please add details about what the pass is supposed to do. tra: Please add details about what the pass is supposed to do.
				// each of them. For example,
				//
				// %A = alloca i32
				// store i32 0, i32* %A ; emits st.u32
				//
				// will be transformed to
				//
				// %A = alloca i32
				// %Local = addrspacecast i32* %A to i32 addrspace(3)*
				// %Shared = addrspacecast i32 addrspace(3)* %A to i32*
				// store i32 0, i32 addrspace(5)* %Shared ; emits st.shared.u32
				//
				// And we will rely on NVPTXInferAddressSpaces to combine the last two
				// instructions.
				//
				// This pass is invoked for -O0 only.
				//
				//===----------------------------------------------------------------------===//

				#include "NVPTX.h"
				#include "NVPTXUtilities.h"
				#include "NVPTXTargetMachine.h"
				#include "llvm/Analysis/ValueTracking.h"
				#include "llvm/IR/Function.h"
				#include "llvm/IR/Instructions.h"
				#include "llvm/IR/Module.h"
				#include "llvm/IR/Type.h"
				#include "llvm/Pass.h"

				using namespace llvm;

				namespace llvm {
				void initializeNVPTXFunctionDataSharingPass(PassRegistry &);
				}

				namespace {
				class NVPTXFunctionDataSharing : public FunctionPass {
				bool runOnFunction(Function &F) override;
				bool runOnKernelFunction(Function &F);
				bool runOnDeviceFunction(Function &F);

				public:
				static char ID; // Pass identification, replacement for typeid
				NVPTXFunctionDataSharing(const NVPTXTargetMachine *TM = nullptr)
				: FunctionPass(ID) {}
				StringRef getPassName() const override {
				return "Function level data sharing pass.";
				}
				};
				} // namespace

				char NVPTXFunctionDataSharing::ID = 1;

				INITIALIZE_PASS(NVPTXFunctionDataSharing, "nvptx-function-data-sharing",
				"Function Data Sharing (NVPTX)", false, false)

				static void markPointerAsShared(Value *Ptr) {
				if (Ptr->getType()->getPointerAddressSpace() == ADDRESS_SPACE_SHARED)
				return;

				// Deciding where to emit the addrspacecast pair.
				// Insert right after Ptr if Ptr is an instruction.
				BasicBlock::iterator InsertPt =
				std::next(cast<Instruction>(Ptr)->getIterator());
				assert(InsertPt != InsertPt->getParent()->end() &&
				"We don't call this function with Ptr being a terminator.");

				auto *PtrInShared = new AddrSpaceCastInst(
				Ptr, PointerType::get(Ptr->getType()->getPointerElementType(),
				ADDRESS_SPACE_SHARED),
				Ptr->getName(), &*InsertPt);
				// Old version
				auto *PtrInGeneric = new AddrSpaceCastInst(PtrInShared, Ptr->getType(),
				Ptr->getName(), &*InsertPt);
				// Replace with PtrInGeneric all uses of Ptr except PtrInShared.
				Ptr->replaceAllUsesWith(PtrInGeneric);
				PtrInShared->setOperand(0, Ptr);
				}

				// =============================================================================
				// Main function for this pass.
				// =============================================================================
				bool NVPTXFunctionDataSharing::runOnKernelFunction(Function &F) {
				bool Modified = false;

				// Skip pass if no data sharing is required.
				if (!F.hasFnAttribute("has-nvptx-shared-depot"))
				return false;

				traUnsubmitted Not Done Reply Inline Actions Nit: `return false` would match the intent better. tra: Nit: `return false` would match the intent better.
				for (auto &B : F) {
				for (auto &I : B) {
				auto *AI = dyn_cast<AllocaInst>(&I);
				if (!AI)
				continue;
				if (AI->getType()->isPointerTy() && ptrIsStored(AI)) {
				markPointerAsShared(AI);
				Modified = true;
				}
				}
				}

				return Modified;
				}

				// Device functions only need to copy byval args into local memory.
				bool NVPTXFunctionDataSharing::runOnDeviceFunction(Function &F) {
				return true;
				}

				bool NVPTXFunctionDataSharing::runOnFunction(Function &F) {
				return isKernelFunction(F) ? runOnKernelFunction(F) : runOnDeviceFunction(F);
				}

				FunctionPass *
				llvm::createNVPTXFunctionDataSharingPass(const NVPTXTargetMachine *TM) {
				return new NVPTXFunctionDataSharing(TM);
				}