1 files changed, 256 insertions, 0 deletions
diff --git a/src/gallium/drivers/radeon/AMDILGlobalManager.h b/src/gallium/drivers/radeon/AMDILGlobalManager.h
new file mode 100644
index 00000000000..1b0361e0174
--- /dev/null
+++ b/src/gallium/drivers/radeon/AMDILGlobalManager.h
@@ -0,0 +1,256 @@
+//===-- 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_