diff options
Diffstat (limited to 'src/gallium/drivers/radeon/AMDILGlobalManager.h')
| -rw-r--r-- | src/gallium/drivers/radeon/AMDILGlobalManager.h | 256 |
1 files changed, 0 insertions, 256 deletions
diff --git a/src/gallium/drivers/radeon/AMDILGlobalManager.h b/src/gallium/drivers/radeon/AMDILGlobalManager.h deleted file mode 100644 index 1b0361e0174..00000000000 --- a/src/gallium/drivers/radeon/AMDILGlobalManager.h +++ /dev/null @@ -1,256 +0,0 @@ -//===-- AMDILGlobalManager.h - TODO: Add brief description -------===// -// -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. -// -// ==-----------------------------------------------------------------------===// -// -// Class that handles parsing and storing global variables that are relevant to -// the compilation of the module. -// -// ==-----------------------------------------------------------------------===// - -#ifndef _AMDILGLOBALMANAGER_H_ -#define _AMDILGLOBALMANAGER_H_ - -#include "AMDIL.h" -#include "llvm/ADT/DenseMap.h" -#include "llvm/ADT/DenseSet.h" -#include "llvm/ADT/SmallSet.h" -#include "llvm/ADT/SmallVector.h" -#include "llvm/ADT/StringMap.h" -#include "llvm/Module.h" -#include "llvm/Support/raw_ostream.h" - -#include <set> -#include <string> - -#define CB_BASE_OFFSET 2 - -namespace llvm { - -class PointerType; -class AMDILKernelManager; -class AMDILSubtarget; -class TypeSymbolTable; -class Argument; -class GlobalValue; -class MachineFunction; - -/// structure that holds information for a single local/region address array -typedef struct _arrayMemRec { - uint32_t vecSize; // size of each vector - uint32_t offset; // offset into the memory section - bool isHW; // flag to specify if HW is used or SW is used - bool isRegion; // flag to specify if GDS is used or not -} arraymem; - -/// Structure that holds information for all local/region address -/// arrays in the kernel -typedef struct _localArgRec { - llvm::SmallVector<arraymem *, DEFAULT_VEC_SLOTS> local; - std::string name; // Kernel Name -} localArg; - -/// structure that holds information about a constant address -/// space pointer that is a kernel argument -typedef struct _constPtrRec { - const Value *base; - uint32_t size; - uint32_t offset; - uint32_t cbNum; // value of 0 means that it does not use hw CB - bool isArray; - bool isArgument; - bool usesHardware; - std::string name; -} constPtr; - -/// Structure that holds information for each kernel argument -typedef struct _kernelArgRec { - uint32_t reqGroupSize[3]; - uint32_t reqRegionSize[3]; - llvm::SmallVector<uint32_t, DEFAULT_VEC_SLOTS> argInfo; - bool mHasRWG; - bool mHasRWR; -} kernelArg; - -/// Structure that holds information for each kernel -typedef struct _kernelRec { - mutable uint32_t curSize; - mutable uint32_t curRSize; - mutable uint32_t curHWSize; - mutable uint32_t curHWRSize; - uint32_t constSize; - kernelArg *sgv; - localArg *lvgv; - llvm::SmallVector<struct _constPtrRec, DEFAULT_VEC_SLOTS> constPtr; - uint32_t constSizes[HW_MAX_NUM_CB]; - llvm::SmallSet<uint32_t, OPENCL_MAX_READ_IMAGES> readOnly; - llvm::SmallSet<uint32_t, OPENCL_MAX_WRITE_IMAGES> writeOnly; - llvm::SmallVector<std::pair<uint32_t, const Constant *>, - DEFAULT_VEC_SLOTS> CPOffsets; -} kernel; - -class AMDILGlobalManager { -public: - AMDILGlobalManager(bool debugMode = false); - ~AMDILGlobalManager(); - - /// Process the given module and parse out the global variable metadata passed - /// down from the frontend-compiler - void processModule(const Module &MF, const AMDILTargetMachine* mTM); - - /// Returns whether the current name is the name of a kernel function or a - /// normal function - bool isKernel(const llvm::StringRef &name) const; - - /// Returns true if the image ID corresponds to a read only image. - bool isReadOnlyImage(const llvm::StringRef &name, uint32_t iID) const; - - /// Returns true if the image ID corresponds to a write only image. - bool isWriteOnlyImage(const llvm::StringRef &name, uint32_t iID) const; - - /// Returns the number of write only images for the kernel. - uint32_t getNumWriteImages(const llvm::StringRef &name) const; - - /// Gets the group size of the kernel for the given dimension. - uint32_t getLocal(const llvm::StringRef &name, uint32_t dim) const; - - /// Gets the region size of the kernel for the given dimension. - uint32_t getRegion(const llvm::StringRef &name, uint32_t dim) const; - - /// Get the Region memory size in 1d for the given function/kernel. - uint32_t getRegionSize(const llvm::StringRef &name) const; - - /// Get the region memory size in 1d for the given function/kernel. - uint32_t getLocalSize(const llvm::StringRef &name) const; - - // Get the max group size in one 1D for the given function/kernel. - uint32_t getMaxGroupSize(const llvm::StringRef &name) const; - - // Get the max region size in one 1D for the given function/kernel. - uint32_t getMaxRegionSize(const llvm::StringRef &name) const; - - /// Get the constant memory size in 1d for the given function/kernel. - uint32_t getConstSize(const llvm::StringRef &name) const; - - /// Get the HW local size in 1d for the given function/kernel We need to - /// seperate SW local and HW local for the case where some local memory is - /// emulated in global and some is using the hardware features. The main - /// problem is that in OpenCL 1.0/1.1 cl_khr_byte_addressable_store allows - /// these actions to happen on all memory spaces, but the hardware can only - /// write byte address stores to UAV and LDS, not GDS or Stack. - uint32_t getHWLocalSize(const llvm::StringRef &name) const; - uint32_t getHWRegionSize(const llvm::StringRef &name) const; - - /// Get the offset of the array for the kernel. - int32_t getArrayOffset(const llvm::StringRef &name) const; - - /// Get the offset of the const memory for the kernel. - int32_t getConstOffset(const llvm::StringRef &name) const; - - /// Get the boolean value if this particular constant uses HW or not. - bool getConstHWBit(const llvm::StringRef &name) const; - - /// Get a reference to the kernel metadata information for the given function - /// name. - const kernel &getKernel(const llvm::StringRef &name) const; - - /// Returns whether a reqd_workgroup_size attribute has been used or not. - bool hasRWG(const llvm::StringRef &name) const; - - /// Returns whether a reqd_workregion_size attribute has been used or not. - bool hasRWR(const llvm::StringRef &name) const; - - - /// Dump the data section to the output stream for the given kernel. - void dumpDataSection(llvm::raw_ostream &O, AMDILKernelManager *km); - - /// Iterate through the constants that are global to the compilation unit. - StringMap<constPtr>::iterator consts_begin(); - StringMap<constPtr>::iterator consts_end(); - - /// Query if the kernel has a byte store. - bool byteStoreExists(llvm::StringRef S) const; - - /// Query if the kernel and argument uses hardware constant memory. - bool usesHWConstant(const kernel &krnl, const llvm::StringRef &arg); - - /// Query if the constant pointer is an argument. - bool isConstPtrArgument(const kernel &krnl, const llvm::StringRef &arg); - - /// Query if the constant pointer is an array that is globally scoped. - bool isConstPtrArray(const kernel &krnl, const llvm::StringRef &arg); - - /// Query the size of the constant pointer. - uint32_t getConstPtrSize(const kernel &krnl, const llvm::StringRef &arg); - - /// Query the offset of the constant pointer. - uint32_t getConstPtrOff(const kernel &krnl, const llvm::StringRef &arg); - - /// Query the constant buffer number for a constant pointer. - uint32_t getConstPtrCB(const kernel &krnl, const llvm::StringRef &arg); - - /// Query the Value* that the constant pointer originates from. - const Value *getConstPtrValue(const kernel &krnl, const llvm::StringRef &arg); - - /// Get the ID of the argument. - int32_t getArgID(const Argument *arg); - - /// Get the unique function ID for the specific function name and create a new - /// unique ID if it is not found. - uint32_t getOrCreateFunctionID(const GlobalValue* func); - uint32_t getOrCreateFunctionID(const std::string& func); - - /// Calculate the offsets of the constant pool for the given kernel and - /// machine function. - void calculateCPOffsets(const MachineFunction *MF, kernel &krnl); - - /// Print the global manager to the output stream. - void print(llvm::raw_ostream& O); - - /// Dump the global manager to the output stream - debug use. - void dump(); - -private: - /// Various functions that parse global value information and store them in - /// the global manager. This approach is used instead of dynamic parsing as it - /// might require more space, but should allow caching of data that gets - /// requested multiple times. - kernelArg parseSGV(const GlobalValue *GV); - localArg parseLVGV(const GlobalValue *GV); - void parseGlobalAnnotate(const GlobalValue *G); - void parseImageAnnotate(const GlobalValue *G); - void parseConstantPtrAnnotate(const GlobalValue *G); - void printConstantValue(const Constant *CAval, - llvm::raw_ostream& O, - bool asByte); - void parseKernelInformation(const Value *V); - void parseAutoArray(const GlobalValue *G, bool isRegion); - void parseConstantPtr(const GlobalValue *G); - void allocateGlobalCB(); - void dumpDataToCB(llvm::raw_ostream &O, AMDILKernelManager *km, uint32_t id); - bool checkConstPtrsUseHW(Module::const_iterator *F); - - llvm::StringMap<arraymem> mArrayMems; - llvm::StringMap<localArg> mLocalArgs; - llvm::StringMap<kernelArg> mKernelArgs; - llvm::StringMap<kernel> mKernels; - llvm::StringMap<constPtr> mConstMems; - llvm::StringMap<uint32_t> mFuncNames; - llvm::DenseMap<const GlobalValue*, uint32_t> mFuncPtrNames; - llvm::DenseMap<uint32_t, llvm::StringRef> mImageNameMap; - std::set<llvm::StringRef> mByteStore; - std::set<llvm::StringRef> mIgnoreStr; - llvm::DenseMap<const Argument *, int32_t> mArgIDMap; - const char *symTab; - const AMDILSubtarget *mSTM; - size_t mOffset; - uint32_t mReservedBuffs; - uint32_t mCurrentCPOffset; - bool mDebugMode; -}; -} // namespace llvm -#endif // __AMDILGLOBALMANAGER_H_ |