CbC/CbC_llvm: lib/Target/NVPTX/NVPTXLowerKernelArgs.cpp annotate

author	Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp>
date	Tue, 13 Oct 2015 17:48:58 +0900
parents
children	7d135dc70f03

rev	line source
95 afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	1 //===-- NVPTXLowerKernelArgs.cpp - Lower kernel arguments -----------------===//
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	2 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	3 // The LLVM Compiler Infrastructure
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	4 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	5 // This file is distributed under the University of Illinois Open Source
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	6 // License. See LICENSE.TXT for details.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	7 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	8 //===----------------------------------------------------------------------===//
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	9 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	10 // Pointer arguments to kernel functions need to be lowered specially.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	11 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	12 // 1. Copy byval struct args to local memory. This is a preparation for handling
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	13 // cases like
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	14 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	15 // kernel void foo(struct A arg, ...)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	16 // {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	17 // struct A *p = &arg;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	18 // ...
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	19 // ... = p->filed1 ... (this is no generic address for .param)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	20 // p->filed2 = ... (this is no write access to .param)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	21 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	22 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	23 // 2. Convert non-byval pointer arguments of CUDA kernels to pointers in the
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	24 // global address space. This allows later optimizations to emit
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	25 // ld.global./st.global. for accessing these pointer arguments. For
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	26 // example,
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	27 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	28 // define void @foo(float* %input) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	29 // %v = load float, float* %input, align 4
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	30 // ...
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	31 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	32 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	33 // becomes
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	34 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	35 // define void @foo(float* %input) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	36 // %input2 = addrspacecast float* %input to float addrspace(1)*
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	37 // %input3 = addrspacecast float addrspace(1)* %input2 to float*
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	38 // %v = load float, float* %input3, align 4
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	39 // ...
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	40 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	41 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	42 // Later, NVPTXFavorNonGenericAddrSpaces will optimize it to
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	43 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	44 // define void @foo(float* %input) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	45 // %input2 = addrspacecast float* %input to float addrspace(1)*
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	46 // %v = load float, float addrspace(1)* %input2, align 4
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	47 // ...
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	48 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	49 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	50 // 3. Convert pointers in a byval kernel parameter to pointers in the global
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	51 // address space. As #2, it allows NVPTX to emit more ld/st.global. E.g.,
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	52 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	53 // struct S {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	54 // int *x;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	55 // int *y;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	56 // };
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	57 // __global__ void foo(S s) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	58 // int *b = s.y;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	59 // // use b
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	60 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	61 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	62 // "b" points to the global address space. In the IR level,
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	63 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	64 // define void @foo({i32, i32}* byval %input) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	65 // %b_ptr = getelementptr {i32, i32}, {i32, i32}* %input, i64 0, i32 1
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	66 // %b = load i32, i32* %b_ptr
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	67 // ; use %b
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	68 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	69 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	70 // becomes
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	71 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	72 // define void @foo({i32, i32}* byval %input) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	73 // %b_ptr = getelementptr {i32, i32}, {i32, i32}* %input, i64 0, i32 1
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	74 // %b = load i32, i32* %b_ptr
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	75 // %b_global = addrspacecast i32* %b to i32 addrspace(1)*
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	76 // %b_generic = addrspacecast i32 addrspace(1)* %b_global to i32*
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	77 // ; use %b_generic
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	78 // }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	79 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	80 // TODO: merge this pass with NVPTXFavorNonGenericAddrSpace so that other passes
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	81 // don't cancel the addrspacecast pair this pass emits.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	82 //===----------------------------------------------------------------------===//
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	83
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	84 #include "NVPTX.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	85 #include "NVPTXUtilities.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	86 #include "NVPTXTargetMachine.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	87 #include "llvm/Analysis/ValueTracking.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	88 #include "llvm/IR/Function.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	89 #include "llvm/IR/Instructions.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	90 #include "llvm/IR/Module.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	91 #include "llvm/IR/Type.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	92 #include "llvm/Pass.h"
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	93
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	94 using namespace llvm;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	95
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	96 namespace llvm {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	97 void initializeNVPTXLowerKernelArgsPass(PassRegistry &);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	98 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	99
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	100 namespace {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	101 class NVPTXLowerKernelArgs : public FunctionPass {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	102 bool runOnFunction(Function &F) override;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	103
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	104 // handle byval parameters
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	105 void handleByValParam(Argument *Arg);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	106 // Knowing Ptr must point to the global address space, this function
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	107 // addrspacecasts Ptr to global and then back to generic. This allows
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	108 // NVPTXFavorNonGenericAddrSpace to fold the global-to-generic cast into
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	109 // loads/stores that appear later.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	110 void markPointerAsGlobal(Value *Ptr);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	111
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	112 public:
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	113 static char ID; // Pass identification, replacement for typeid
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	114 NVPTXLowerKernelArgs(const NVPTXTargetMachine *TM = nullptr)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	115 : FunctionPass(ID), TM(TM) {}
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	116 const char *getPassName() const override {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	117 return "Lower pointer arguments of CUDA kernels";
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	118 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	119
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	120 private:
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	121 const NVPTXTargetMachine *TM;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	122 };
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	123 } // namespace
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	124
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	125 char NVPTXLowerKernelArgs::ID = 1;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	126
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	127 INITIALIZE_PASS(NVPTXLowerKernelArgs, "nvptx-lower-kernel-args",
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	128 "Lower kernel arguments (NVPTX)", false, false)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	129
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	130 // =============================================================================
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	131 // If the function had a byval struct ptr arg, say foo(%struct.x *byval %d),
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	132 // then add the following instructions to the first basic block:
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	133 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	134 // %temp = alloca %struct.x, align 8
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	135 // %tempd = addrspacecast %struct.x* %d to %struct.x addrspace(101)*
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	136 // %tv = load %struct.x addrspace(101)* %tempd
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	137 // store %struct.x %tv, %struct.x* %temp, align 8
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	138 //
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	139 // The above code allocates some space in the stack and copies the incoming
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	140 // struct from param space to local space.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	141 // Then replace all occurrences of %d by %temp.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	142 // =============================================================================
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	143 void NVPTXLowerKernelArgs::handleByValParam(Argument *Arg) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	144 Function *Func = Arg->getParent();
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	145 Instruction *FirstInst = &(Func->getEntryBlock().front());
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	146 PointerType *PType = dyn_cast<PointerType>(Arg->getType());
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	147
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	148 assert(PType && "Expecting pointer type in handleByValParam");
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	149
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	150 Type *StructType = PType->getElementType();
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	151 AllocaInst *AllocA = new AllocaInst(StructType, Arg->getName(), FirstInst);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	152 // Set the alignment to alignment of the byval parameter. This is because,
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	153 // later load/stores assume that alignment, and we are going to replace
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	154 // the use of the byval parameter with this alloca instruction.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	155 AllocA->setAlignment(Func->getParamAlignment(Arg->getArgNo() + 1));
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	156 Arg->replaceAllUsesWith(AllocA);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	157
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	158 Value *ArgInParam = new AddrSpaceCastInst(
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	159 Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(),
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	160 FirstInst);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	161 LoadInst *LI = new LoadInst(ArgInParam, Arg->getName(), FirstInst);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	162 new StoreInst(LI, AllocA, FirstInst);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	163 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	164
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	165 void NVPTXLowerKernelArgs::markPointerAsGlobal(Value *Ptr) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	166 if (Ptr->getType()->getPointerAddressSpace() == ADDRESS_SPACE_GLOBAL)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	167 return;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	168
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	169 // Deciding where to emit the addrspacecast pair.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	170 BasicBlock::iterator InsertPt;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	171 if (Argument *Arg = dyn_cast<Argument>(Ptr)) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	172 // Insert at the functon entry if Ptr is an argument.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	173 InsertPt = Arg->getParent()->getEntryBlock().begin();
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	174 } else {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	175 // Insert right after Ptr if Ptr is an instruction.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	176 InsertPt = cast<Instruction>(Ptr);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	177 ++InsertPt;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	178 assert(InsertPt != InsertPt->getParent()->end() &&
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	179 "We don't call this function with Ptr being a terminator.");
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	180 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	181
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	182 Instruction *PtrInGlobal = new AddrSpaceCastInst(
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	183 Ptr, PointerType::get(Ptr->getType()->getPointerElementType(),
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	184 ADDRESS_SPACE_GLOBAL),
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	185 Ptr->getName(), InsertPt);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	186 Value *PtrInGeneric = new AddrSpaceCastInst(PtrInGlobal, Ptr->getType(),
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	187 Ptr->getName(), InsertPt);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	188 // Replace with PtrInGeneric all uses of Ptr except PtrInGlobal.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	189 Ptr->replaceAllUsesWith(PtrInGeneric);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	190 PtrInGlobal->setOperand(0, Ptr);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	191 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	192
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	193 // =============================================================================
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	194 // Main function for this pass.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	195 // =============================================================================
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	196 bool NVPTXLowerKernelArgs::runOnFunction(Function &F) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	197 // Skip non-kernels. See the comments at the top of this file.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	198 if (!isKernelFunction(F))
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	199 return false;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	200
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	201 if (TM && TM->getDrvInterface() == NVPTX::CUDA) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	202 // Mark pointers in byval structs as global.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	203 for (auto &B : F) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	204 for (auto &I : B) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	205 if (LoadInst *LI = dyn_cast<LoadInst>(&I)) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	206 if (LI->getType()->isPointerTy()) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	207 Value *UO = GetUnderlyingObject(LI->getPointerOperand(),
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	208 F.getParent()->getDataLayout());
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	209 if (Argument *Arg = dyn_cast<Argument>(UO)) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	210 if (Arg->hasByValAttr()) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	211 // LI is a load from a pointer within a byval kernel parameter.
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	212 markPointerAsGlobal(LI);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	213 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	214 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	215 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	216 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	217 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	218 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	219 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	220
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	221 for (Argument &Arg : F.args()) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	222 if (Arg.getType()->isPointerTy()) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	223 if (Arg.hasByValAttr())
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	224 handleByValParam(&Arg);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	225 else if (TM && TM->getDrvInterface() == NVPTX::CUDA)
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	226 markPointerAsGlobal(&Arg);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	227 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	228 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	229 return true;
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	230 }
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	231
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	232 FunctionPass *
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	233 llvm::createNVPTXLowerKernelArgsPass(const NVPTXTargetMachine *TM) {
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	234 return new NVPTXLowerKernelArgs(TM);
afa8332a0e37 LLVM 3.8 Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp> parents: diff changeset	235 }

95

afa8332a0e37 LLVM 3.8