src/gallium/drivers/radeon/AMDILGlobalManager.h

   1 //===-- AMDILGlobalManager.h - TODO: Add brief description -------===//
   2 //
   3 //                     The LLVM Compiler Infrastructure
   4 //
   5 // This file is distributed under the University of Illinois Open Source
   6 // License. See LICENSE.TXT for details.
   7 //
   8 // ==-----------------------------------------------------------------------===//
   9 //
  10 // Class that handles parsing and storing global variables that are relevant to
  11 // the compilation of the module.
  12 //
  13 // ==-----------------------------------------------------------------------===//
  14
  15 #ifndef _AMDILGLOBALMANAGER_H_
  16 #define _AMDILGLOBALMANAGER_H_
  17
  18 #include "AMDIL.h"
  19 #include "llvm/ADT/DenseMap.h"
  20 #include "llvm/ADT/DenseSet.h"
  21 #include "llvm/ADT/SmallSet.h"
  22 #include "llvm/ADT/SmallVector.h"
  23 #include "llvm/ADT/StringMap.h"
  24 #include "llvm/Module.h"
  25 #include "llvm/Support/raw_ostream.h"
  26
  27 #include <set>
  28 #include <string>
  29
  30 #define CB_BASE_OFFSET 2
  31
  32 namespace llvm {
  33
  34 class PointerType;
  35 class AMDILKernelManager;
  36 class AMDILSubtarget;
  37 class TypeSymbolTable;
  38 class Argument;
  39 class GlobalValue;
  40 class MachineFunction;
  41
  42 /// structure that holds information for a single local/region address array
  43 typedef struct _arrayMemRec {
  44   uint32_t vecSize; // size of each vector
  45   uint32_t offset;  // offset into the memory section
  46   bool isHW;        // flag to specify if HW is used or SW is used
  47   bool isRegion;    // flag to specify if GDS is used or not
  48 } arraymem;
  49
  50 /// Structure that holds information for all local/region address
  51 /// arrays in the kernel
  52 typedef struct _localArgRec {
  53   llvm::SmallVector<arraymem *, DEFAULT_VEC_SLOTS> local;
  54   std::string name; // Kernel Name
  55 } localArg;
  56
  57 /// structure that holds information about a constant address
  58 /// space pointer that is a kernel argument
  59 typedef struct _constPtrRec {
  60   const Value *base;
  61   uint32_t size;
  62   uint32_t offset;
  63   uint32_t cbNum; // value of 0 means that it does not use hw CB
  64   bool isArray;
  65   bool isArgument;
  66   bool usesHardware;
  67   std::string name;
  68 } constPtr;
  69
  70 /// Structure that holds information for each kernel argument
  71 typedef struct _kernelArgRec {
  72   uint32_t reqGroupSize[3];
  73   uint32_t reqRegionSize[3];
  74   llvm::SmallVector<uint32_t, DEFAULT_VEC_SLOTS> argInfo;
  75   bool mHasRWG;
  76   bool mHasRWR;
  77 } kernelArg;
  78
  79 /// Structure that holds information for each kernel
  80 typedef struct _kernelRec {
  81   mutable uint32_t curSize;
  82   mutable uint32_t curRSize;
  83   mutable uint32_t curHWSize;
  84   mutable uint32_t curHWRSize;
  85   uint32_t constSize;
  86   kernelArg *sgv;
  87   localArg *lvgv;
  88   llvm::SmallVector<struct _constPtrRec, DEFAULT_VEC_SLOTS> constPtr;
  89   uint32_t constSizes[HW_MAX_NUM_CB];
  90   llvm::SmallSet<uint32_t, OPENCL_MAX_READ_IMAGES> readOnly;
  91   llvm::SmallSet<uint32_t, OPENCL_MAX_WRITE_IMAGES> writeOnly;
  92   llvm::SmallVector<std::pair<uint32_t, const Constant *>,
  93     DEFAULT_VEC_SLOTS> CPOffsets;
  94 } kernel;
  95
  96 class AMDILGlobalManager {
  97 public:
  98   AMDILGlobalManager(bool debugMode = false);
  99   ~AMDILGlobalManager();
 100
 101   /// Process the given module and parse out the global variable metadata passed
 102   /// down from the frontend-compiler
 103   void processModule(const Module &MF, const AMDILTargetMachine* mTM);
 104
 105   /// Returns whether the current name is the name of a kernel function or a
 106   /// normal function
 107   bool isKernel(const llvm::StringRef &name) const;
 108
 109   /// Returns true if the image ID corresponds to a read only image.
 110   bool isReadOnlyImage(const llvm::StringRef &name, uint32_t iID) const;
 111
 112   /// Returns true if the image ID corresponds to a write only image.
 113   bool isWriteOnlyImage(const llvm::StringRef &name, uint32_t iID) const;
 114
 115   /// Returns the number of write only images for the kernel.
 116   uint32_t getNumWriteImages(const llvm::StringRef &name) const;
 117
 118   /// Gets the group size of the kernel for the given dimension.
 119   uint32_t getLocal(const llvm::StringRef &name, uint32_t dim) const;
 120
 121   /// Gets the region size of the kernel for the given dimension.
 122   uint32_t getRegion(const llvm::StringRef &name, uint32_t dim) const;
 123
 124   /// Get the Region memory size in 1d for the given function/kernel.
 125   uint32_t getRegionSize(const llvm::StringRef &name) const;
 126
 127   /// Get the region memory size in 1d for the given function/kernel.
 128   uint32_t getLocalSize(const llvm::StringRef &name) const;
 129
 130   // Get the max group size in one 1D for the given function/kernel.
 131   uint32_t getMaxGroupSize(const llvm::StringRef &name) const;
 132
 133   // Get the max region size in one 1D for the given function/kernel.
 134   uint32_t getMaxRegionSize(const llvm::StringRef &name) const;
 135
 136   /// Get the constant memory size in 1d for the given function/kernel.
 137   uint32_t getConstSize(const llvm::StringRef &name) const;
 138
 139   /// Get the HW local size in 1d for the given function/kernel We need to
 140   /// seperate SW local and HW local for the case where some local memory is
 141   /// emulated in global and some is using the hardware features. The main
 142   /// problem is that in OpenCL 1.0/1.1 cl_khr_byte_addressable_store allows
 143   /// these actions to happen on all memory spaces, but the hardware can only
 144   /// write byte address stores to UAV and LDS, not GDS or Stack.
 145   uint32_t getHWLocalSize(const llvm::StringRef &name) const;
 146   uint32_t getHWRegionSize(const llvm::StringRef &name) const;
 147
 148   /// Get the offset of the array for the kernel.
 149   int32_t getArrayOffset(const llvm::StringRef &name) const;
 150
 151   /// Get the offset of the const memory for the kernel.
 152   int32_t getConstOffset(const llvm::StringRef &name) const;
 153
 154   /// Get the boolean value if this particular constant uses HW or not.
 155   bool getConstHWBit(const llvm::StringRef &name) const;
 156
 157   /// Get a reference to the kernel metadata information for the given function
 158   /// name.
 159   const kernel &getKernel(const llvm::StringRef &name) const;
 160
 161   /// Returns whether a reqd_workgroup_size attribute has been used or not.
 162   bool hasRWG(const llvm::StringRef &name) const;
 163
 164   /// Returns whether a reqd_workregion_size attribute has been used or not.
 165   bool hasRWR(const llvm::StringRef &name) const;
 166
 167
 168   /// Dump the data section to the output stream for the given kernel.
 169   void dumpDataSection(llvm::raw_ostream &O, AMDILKernelManager *km);
 170
 171   /// Iterate through the constants that are global to the compilation unit.
 172   StringMap<constPtr>::iterator consts_begin();
 173   StringMap<constPtr>::iterator consts_end();
 174
 175   /// Query if the kernel has a byte store.
 176   bool byteStoreExists(llvm::StringRef S) const;
 177
 178   /// Query if the kernel and argument uses hardware constant memory.
 179   bool usesHWConstant(const kernel &krnl, const llvm::StringRef &arg);
 180
 181   /// Query if the constant pointer is an argument.
 182   bool isConstPtrArgument(const kernel &krnl, const llvm::StringRef &arg);
 183
 184   /// Query if the constant pointer is an array that is globally scoped.
 185   bool isConstPtrArray(const kernel &krnl, const llvm::StringRef &arg);
 186
 187   /// Query the size of the constant pointer.
 188   uint32_t getConstPtrSize(const kernel &krnl, const llvm::StringRef &arg);
 189
 190   /// Query the offset of the constant pointer.
 191   uint32_t getConstPtrOff(const kernel &krnl, const llvm::StringRef &arg);
 192
 193   /// Query the constant buffer number for a constant pointer.
 194   uint32_t getConstPtrCB(const kernel &krnl, const llvm::StringRef &arg);
 195
 196   /// Query the Value* that the constant pointer originates from.
 197   const Value *getConstPtrValue(const kernel &krnl, const llvm::StringRef &arg);
 198
 199   /// Get the ID of the argument.
 200   int32_t getArgID(const Argument *arg);
 201
 202   /// Get the unique function ID for the specific function name and create a new
 203   /// unique ID if it is not found.
 204   uint32_t getOrCreateFunctionID(const GlobalValue* func);
 205   uint32_t getOrCreateFunctionID(const std::string& func);
 206
 207   /// Calculate the offsets of the constant pool for the given kernel and
 208   /// machine function.
 209   void calculateCPOffsets(const MachineFunction *MF, kernel &krnl);
 210
 211   /// Print the global manager to the output stream.
 212   void print(llvm::raw_ostream& O);
 213
 214   /// Dump the global manager to the output stream - debug use.
 215   void dump();
 216
 217 private:
 218   /// Various functions that parse global value information and store them in
 219   /// the global manager. This approach is used instead of dynamic parsing as it
 220   /// might require more space, but should allow caching of data that gets
 221   /// requested multiple times.
 222   kernelArg parseSGV(const GlobalValue *GV);
 223   localArg  parseLVGV(const GlobalValue *GV);
 224   void parseGlobalAnnotate(const GlobalValue *G);
 225   void parseImageAnnotate(const GlobalValue *G);
 226   void parseConstantPtrAnnotate(const GlobalValue *G);
 227   void printConstantValue(const Constant *CAval,
 228                           llvm::raw_ostream& O,
 229                           bool asByte);
 230   void parseKernelInformation(const Value *V);
 231   void parseAutoArray(const GlobalValue *G, bool isRegion);
 232   void parseConstantPtr(const GlobalValue *G);
 233   void allocateGlobalCB();
 234   void dumpDataToCB(llvm::raw_ostream &O, AMDILKernelManager *km, uint32_t id);
 235   bool checkConstPtrsUseHW(Module::const_iterator *F);
 236
 237   llvm::StringMap<arraymem> mArrayMems;
 238   llvm::StringMap<localArg> mLocalArgs;
 239   llvm::StringMap<kernelArg> mKernelArgs;
 240   llvm::StringMap<kernel> mKernels;
 241   llvm::StringMap<constPtr> mConstMems;
 242   llvm::StringMap<uint32_t> mFuncNames;
 243   llvm::DenseMap<const GlobalValue*, uint32_t> mFuncPtrNames;
 244   llvm::DenseMap<uint32_t, llvm::StringRef> mImageNameMap;
 245   std::set<llvm::StringRef> mByteStore;
 246   std::set<llvm::StringRef> mIgnoreStr;
 247   llvm::DenseMap<const Argument *, int32_t> mArgIDMap;
 248   const char *symTab;
 249   const AMDILSubtarget *mSTM;
 250   size_t mOffset;
 251   uint32_t mReservedBuffs;
 252   uint32_t mCurrentCPOffset;
 253   bool mDebugMode;
 254 };
 255 } // namespace llvm
 256 #endif // __AMDILGLOBALMANAGER_H_