radeon/llvm: Remove GlobalManager and KernelManager

11 files changed, 23 insertions, 3275 deletions

diff --git a/src/gallium/drivers/radeon/AMDGPUTargetMachine.cpp b/src/gallium/drivers/radeon/AMDGPUTargetMachine.cpp
index b006f84629e..f8e1cd95dc9 100644
--- a/src/gallium/drivers/radeon/AMDGPUTargetMachine.cpp
+++ b/src/gallium/drivers/radeon/AMDGPUTargetMachine.cpp
@@ -13,8 +13,6 @@
 
 #include "AMDGPUTargetMachine.h"
 #include "AMDGPU.h"
-#include "AMDILGlobalManager.h"
-#include "AMDILKernelManager.h"
 #include "AMDILTargetMachine.h"
 #include "R600ISelLowering.h"
 #include "R600InstrInfo.h"
@@ -44,16 +42,9 @@ AMDGPUTargetMachine::AMDGPUTargetMachine(const Target &T, StringRef TT,
 :
   AMDILTargetMachine(T, TT, CPU, FS, Options, RM, CM, OptLevel),
   Subtarget(TT, CPU, FS),
-  mGM(new AMDILGlobalManager(0 /* Debug mode */)),
-  mKM(new AMDILKernelManager(this, mGM)),
   mDump(false)
 
 {
-  /* XXX: Add these two initializations to fix a segfault, not sure if this
-   * is correct.  These are normally initialized in the AsmPrinter, but AMDGPU
-   * does not use the asm printer */
-  Subtarget.setGlobalManager(mGM);
-  Subtarget.setKernelManager(mKM);
   /* TLInfo uses InstrInfo so it must be initialized after. */
   if (Subtarget.device()->getGeneration() <= AMDILDeviceInfo::HD6XXX) {
     InstrInfo = new R600InstrInfo(*this);
@@ -66,8 +57,6 @@ AMDGPUTargetMachine::AMDGPUTargetMachine(const Target &T, StringRef TT,
 
 AMDGPUTargetMachine::~AMDGPUTargetMachine()
 {
-    delete mGM;
-    delete mKM;
 }
 
 bool AMDGPUTargetMachine::addPassesToEmitFile(PassManagerBase &PM,
diff --git a/src/gallium/drivers/radeon/AMDILGlobalManager.cpp b/src/gallium/drivers/radeon/AMDILGlobalManager.cpp
deleted file mode 100644
index eafd36eaa4e..00000000000
--- a/src/gallium/drivers/radeon/AMDILGlobalManager.cpp
+++ /dev/null
@@ -1,1353 +0,0 @@
-//===-- AMDILGlobalManager.cpp - 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.
-//
-//==-----------------------------------------------------------------------===//
-#include "AMDILGlobalManager.h"
-#include "AMDILDevices.h"
-#include "AMDILKernelManager.h"
-#include "AMDILSubtarget.h"
-
-#include "AMDILAlgorithms.tpp"
-#include "AMDILGlobalManager.h"
-#include "AMDILDevices.h"
-#include "AMDILKernelManager.h"
-#include "AMDILSubtarget.h"
-#include "AMDILUtilityFunctions.h"
-#include "llvm/CodeGen/MachineConstantPool.h"
-#include "llvm/Constants.h"
-#include "llvm/DerivedTypes.h"
-#include "llvm/Instructions.h"
-#include "llvm/Support/FormattedStream.h"
-
-#include <cstdio>
-
-using namespace llvm;
-
-AMDILGlobalManager::AMDILGlobalManager(bool debugMode) {
-  mOffset = 0;
-  mReservedBuffs = 0;
-  symTab = NULL;
-  mCurrentCPOffset = 0;
-  mDebugMode = debugMode;
-}
-
-AMDILGlobalManager::~AMDILGlobalManager() {
-}
-
-void AMDILGlobalManager::print(llvm::raw_ostream &O) {
-  if (!mDebugMode) {
-    return;
-  }
-  O << ";AMDIL Global Manager State Dump:\n";
-  O << ";\tSubtarget: " << mSTM << "\tSymbol Table: " << symTab
-    << "\n";
-  O << ";\tConstant Offset: " << mOffset << "\tCP Offset: "
-    << mCurrentCPOffset << "\tReserved Buffers: " << mReservedBuffs
-    << "\n";
-  if (!mImageNameMap.empty()) {
-    llvm::DenseMap<uint32_t, llvm::StringRef>::iterator imb, ime;
-    O << ";\tGlobal Image Mapping: \n";
-    for (imb = mImageNameMap.begin(), ime = mImageNameMap.end(); imb != ime;
-         ++imb) {
-      O << ";\t\tImage ID: " << imb->first << "\tName: "
-        << imb->second << "\n";
-    }
-  }
-  std::set<llvm::StringRef>::iterator sb, se;
-  if (!mByteStore.empty()) {
-    O << ";Byte Store Kernels: \n";
-    for (sb = mByteStore.begin(), se = mByteStore.end(); sb != se; ++sb) {
-      O << ";\t\t" << *sb << "\n";
-    }
-  }
-  if (!mIgnoreStr.empty()) {
-    O << ";\tIgnored Data Strings: \n";
-    for (sb = mIgnoreStr.begin(), se = mIgnoreStr.end(); sb != se; ++sb) {
-      O << ";\t\t" << *sb << "\n";
-    }
-  }
-}
-
-void AMDILGlobalManager::dump() {
-  print(errs());
-}
-
-static const constPtr *getConstPtr(const kernel &krnl, const std::string &arg) {
-  llvm::SmallVector<constPtr, DEFAULT_VEC_SLOTS>::const_iterator begin, end;
-  for (begin = krnl.constPtr.begin(), end = krnl.constPtr.end();
-       begin != end; ++begin) {
-    if (!strcmp(begin->name.data(),arg.c_str())) {
-      return &(*begin);
-    }
-  }
-  return NULL;
-}
-#if 0
-static bool structContainsSub32bitType(const StructType *ST) {
-  StructType::element_iterator eib, eie;
-  for (eib = ST->element_begin(), eie = ST->element_end(); eib != eie; ++eib) {
-    Type *ptr = *eib;
-    uint32_t size = (uint32_t)GET_SCALAR_SIZE(ptr);
-    if (!size) {
-      if (const StructType *ST = dyn_cast<StructType>(ptr)) {
-        if (structContainsSub32bitType(ST)) {
-          return true;
-        }
-      }
-    } else if (size < 32) {
-      return true;
-    }
-  }
-  return false;
-}
-#endif
-
-void AMDILGlobalManager::processModule(const Module &M,
-                                       const AMDILTargetMachine *mTM)
-{
-  Module::const_global_iterator GI;
-  Module::const_global_iterator GE;
-  symTab = "NoSymTab";
-  mSTM = mTM->getSubtargetImpl();
-  for (GI = M.global_begin(), GE = M.global_end(); GI != GE; ++GI) {
-    const GlobalValue *GV = GI;
-    if (mDebugMode) {
-      GV->dump();
-      errs() << "\n";
-    }
-    llvm::StringRef GVName = GV->getName();
-    const char *name = GVName.data();
-    if (!strncmp(name, "sgv", 3)) {
-      mKernelArgs[GVName] = parseSGV(GV);
-    } else if (!strncmp(name, "fgv", 3)) {
-      // we can ignore this since we don't care about the filename
-      // string
-    } else if (!strncmp(name, "lvgv", 4)) {
-      mLocalArgs[GVName] = parseLVGV(GV);
-    } else if (!strncmp(name, "llvm.image.annotations", 22)) {
-      if (strstr(name, "__OpenCL")
-          && strstr(name, "_kernel")) {
-        // we only want to parse the image information if the
-        // image is a kernel, we might have to parse out the
-        // information if a function is found that is not
-        // inlined.
-        parseImageAnnotate(GV);
-      }
-    } else if (!strncmp(name, "llvm.global.annotations", 23)) {
-      parseGlobalAnnotate(GV);
-    } else if (!strncmp(name, "llvm.constpointer.annotations", 29)) {
-      if (strstr(name, "__OpenCL")
-          && strstr(name, "_kernel")) {
-        // we only want to parse constant pointer information
-        // if it is a kernel
-        parseConstantPtrAnnotate(GV);
-      }
-    } else if (!strncmp(name, "llvm.readonlypointer.annotations", 32)) {
-      // These are skipped as we handle them later in AMDILPointerManager.cpp
-    } else if (GV->getType()->getAddressSpace() == 3) { // *** Match cl_kernel.h local AS #
-      parseAutoArray(GV, false);
-    } else if (strstr(name, "clregion")) {
-      parseAutoArray(GV, true);
-    } else if (!GV->use_empty()
-               && mIgnoreStr.find(GVName) == mIgnoreStr.end()) {
-      parseConstantPtr(GV);
-    }
-  }
-  allocateGlobalCB();
-
-  safeForEach(M.begin(), M.end(),
-      std::bind1st(
-        std::mem_fun(&AMDILGlobalManager::checkConstPtrsUseHW),
-        this));
-}
-
-void AMDILGlobalManager::allocateGlobalCB(void) {
-  uint32_t maxCBSize = mSTM->device()->getMaxCBSize();
-  uint32_t offset = 0;
-  uint32_t curCB = 0;
-  uint32_t swoffset = 0;
-  for (StringMap<constPtr>::iterator cpb = mConstMems.begin(),
-       cpe = mConstMems.end(); cpb != cpe; ++cpb) {
-    bool constHW = mSTM->device()->usesHardware(AMDILDeviceInfo::ConstantMem);
-    cpb->second.usesHardware = false;
-    if (constHW) {
-      // If we have a limit on the max CB Size, then we need to make sure that
-      // the constant sizes fall within the limits.
-      if (cpb->second.size <= maxCBSize) {
-        if (offset + cpb->second.size > maxCBSize) {
-          offset = 0;
-          curCB++;
-        }
-        if (curCB < mSTM->device()->getMaxNumCBs()) {
-          cpb->second.cbNum = curCB + CB_BASE_OFFSET;
-          cpb->second.offset = offset;
-          offset += (cpb->second.size + 15) & (~15);
-          cpb->second.usesHardware = true;
-          continue;
-        }
-      }
-    }
-    cpb->second.cbNum = 0;
-    cpb->second.offset = swoffset;
-    swoffset += (cpb->second.size + 15) & (~15);
-  }
-  if (!mConstMems.empty()) {
-    mReservedBuffs = curCB + 1;
-  }
-}
-
-bool AMDILGlobalManager::checkConstPtrsUseHW(llvm::Module::const_iterator *FCI)
-{
-  Function::const_arg_iterator AI, AE;
-  const Function *func = *FCI;
-  std::string name = func->getName();
-  if (!strstr(name.c_str(), "__OpenCL")
-      || !strstr(name.c_str(), "_kernel")) {
-    return false;
-  }
-  kernel &krnl =  mKernels[name];
-  if (mSTM->device()->usesHardware(AMDILDeviceInfo::ConstantMem)) {
-    for (AI = func->arg_begin(), AE = func->arg_end();
-         AI != AE; ++AI) {
-      const Argument *Arg = &(*AI);
-      const PointerType *P = dyn_cast<PointerType>(Arg->getType());
-      if (!P) {
-        continue;
-      }
-      if (P->getAddressSpace() != AMDILAS::CONSTANT_ADDRESS) {
-        continue;
-      }
-      const constPtr *ptr = getConstPtr(krnl, Arg->getName());
-      if (ptr) {
-        continue;
-      }
-      constPtr constAttr;
-      constAttr.name = Arg->getName();
-      constAttr.size = this->mSTM->device()->getMaxCBSize();
-      constAttr.base = Arg;
-      constAttr.isArgument = true;
-      constAttr.isArray = false;
-      constAttr.offset = 0;
-      constAttr.usesHardware =
-        mSTM->device()->usesHardware(AMDILDeviceInfo::ConstantMem);
-      if (constAttr.usesHardware) {
-        constAttr.cbNum = krnl.constPtr.size() + 2;
-      } else {
-        constAttr.cbNum = 0;
-      }
-      krnl.constPtr.push_back(constAttr);
-    }
-  }
-  // Now lets make sure that only the N largest buffers
-  // get allocated in hardware if we have too many buffers
-  uint32_t numPtrs = krnl.constPtr.size();
-  if (numPtrs > (this->mSTM->device()->getMaxNumCBs() - mReservedBuffs)) {
-    // TODO: Change this routine so it sorts
-    // constPtr instead of pulling the sizes out
-    // and then grab the N largest and disable the rest
-    llvm::SmallVector<uint32_t, 16> sizes;
-    for (uint32_t x = 0; x < numPtrs; ++x) {
-      sizes.push_back(krnl.constPtr[x].size);
-    }
-    std::sort(sizes.begin(), sizes.end());
-    uint32_t numToDisable = numPtrs - (mSTM->device()->getMaxNumCBs() -
-                                       mReservedBuffs);
-    uint32_t safeSize = sizes[numToDisable-1];
-    for (uint32_t x = 0; x < numPtrs && numToDisable; ++x) {
-      if (krnl.constPtr[x].size <= safeSize) {
-        krnl.constPtr[x].usesHardware = false;
-        --numToDisable;
-      }
-    }
-  }
-  // Renumber all of the valid CB's so that
-  // they are linear increase
-  uint32_t CBid = 2 + mReservedBuffs;
-  for (uint32_t x = 0; x < numPtrs; ++x) {
-    if (krnl.constPtr[x].usesHardware) {
-      krnl.constPtr[x].cbNum = CBid++;
-    }
-  }
-  for (StringMap<constPtr>::iterator cpb = mConstMems.begin(),
-       cpe = mConstMems.end(); cpb != cpe; ++cpb) {
-    if (cpb->second.usesHardware) {
-      krnl.constPtr.push_back(cpb->second);
-    }
-  }
-  for (uint32_t x = 0; x < krnl.constPtr.size(); ++x) {
-    constPtr &c = krnl.constPtr[x];
-    uint32_t cbNum = c.cbNum - CB_BASE_OFFSET;
-    if (cbNum < HW_MAX_NUM_CB && c.cbNum >= CB_BASE_OFFSET) {
-      if ((c.size + c.offset) > krnl.constSizes[cbNum]) {
-        krnl.constSizes[cbNum] =
-          ((c.size + c.offset) + 15) & ~15;
-      }
-    } else {
-      krnl.constPtr[x].usesHardware = false;
-    }
-  }
-  return false;
-}
-
-int32_t AMDILGlobalManager::getArrayOffset(const llvm::StringRef &a) const {
-  StringMap<arraymem>::const_iterator iter = mArrayMems.find(a);
-  if (iter != mArrayMems.end()) {
-    return iter->second.offset;
-  } else {
-    return -1;
-  }
-}
-
-int32_t AMDILGlobalManager::getConstOffset(const llvm::StringRef &a) const {
-  StringMap<constPtr>::const_iterator iter = mConstMems.find(a);
-  if (iter != mConstMems.end()) {
-    return iter->second.offset;
-  } else {
-    return -1;
-  }
-}
-
-bool AMDILGlobalManager::getConstHWBit(const llvm::StringRef &name) const {
-  StringMap<constPtr>::const_iterator iter = mConstMems.find(name);
-  if (iter != mConstMems.end()) {
-    return iter->second.usesHardware;
-  } else {
-    return false;
-  }
-}
-
-// As of right now we only care about the required group size
-// so we can skip the variable encoding
-kernelArg AMDILGlobalManager::parseSGV(const GlobalValue *G) {
-  kernelArg nArg;
-  const GlobalVariable *GV = dyn_cast<GlobalVariable>(G);
-  memset(&nArg, 0, sizeof(nArg));
-  for (int x = 0; x < 3; ++x) {
-    nArg.reqGroupSize[x] = mSTM->getDefaultSize(x);
-    nArg.reqRegionSize[x] = mSTM->getDefaultSize(x);
-  }
-  if (!GV || !GV->hasInitializer()) {
-    return nArg;
-  }
-  const Constant *CV = GV->getInitializer();
-  const ConstantDataArray *CA =dyn_cast_or_null<ConstantDataArray>(CV);
-
-  if (!CA || !CA->isString()) {
-    return nArg;
-  }
-  std::string init = CA->getAsString();
-  size_t pos = init.find("RWG");
-  if (pos != llvm::StringRef::npos) {
-    pos += 3;
-    std::string LWS = init.substr(pos, init.length() - pos);
-    const char *lws = LWS.c_str();
-    sscanf(lws, "%u,%u,%u", &(nArg.reqGroupSize[0]),
-           &(nArg.reqGroupSize[1]),
-           &(nArg.reqGroupSize[2]));
-    nArg.mHasRWG = true;
-  }
-  pos = init.find("RWR");
-  if (pos != llvm::StringRef::npos) {
-    pos += 3;
-    std::string LWS = init.substr(pos, init.length() - pos);
-    const char *lws = LWS.c_str();
-    sscanf(lws, "%u,%u,%u", &(nArg.reqRegionSize[0]),
-           &(nArg.reqRegionSize[1]),
-           &(nArg.reqRegionSize[2]));
-    nArg.mHasRWR = true;
-  }
-  return nArg;
-}
-
-localArg AMDILGlobalManager::parseLVGV(const GlobalValue *G) {
-  localArg nArg;
-  const GlobalVariable *GV = dyn_cast<GlobalVariable>(G);
-  nArg.name = "";
-  if (!GV || !GV->hasInitializer()) {
-    return nArg;
-  }
-  const ConstantArray *CA =
-    dyn_cast_or_null<ConstantArray>(GV->getInitializer());
-  if (!CA) {
-    return nArg;
-  }
-  for (size_t x = 0, y = CA->getNumOperands(); x < y; ++x) {
-    const Value *local = CA->getOperand(x);
-    const ConstantExpr *CE = dyn_cast_or_null<ConstantExpr>(local);
-    if (!CE || !CE->getNumOperands()) {
-      continue;
-    }
-    nArg.name = (*(CE->op_begin()))->getName();
-    if (mArrayMems.find(nArg.name) != mArrayMems.end()) {
-      nArg.local.push_back(&(mArrayMems[nArg.name]));
-    }
-  }
-  return nArg;
-}
-
-void AMDILGlobalManager::parseConstantPtrAnnotate(const GlobalValue *G) {
-  const GlobalVariable *GV = dyn_cast_or_null<GlobalVariable>(G);
-  const ConstantArray *CA =
-    dyn_cast_or_null<ConstantArray>(GV->getInitializer());
-  if (!CA) {
-    return;
-  }
-  uint32_t numOps = CA->getNumOperands();
-  for (uint32_t x = 0; x < numOps; ++x) {
-    const Value *V = CA->getOperand(x);
-    const ConstantStruct *CS = dyn_cast_or_null<ConstantStruct>(V);
-    if (!CS) {
-      continue;
-    }
-    assert(CS->getNumOperands() == 2 && "There can only be 2"
-           " fields, a name and size");
-    const ConstantExpr *nameField = dyn_cast<ConstantExpr>(CS->getOperand(0));
-    const ConstantInt *sizeField = dyn_cast<ConstantInt>(CS->getOperand(1));
-    assert(nameField && "There must be a constant name field");
-    assert(sizeField && "There must be a constant size field");
-    const GlobalVariable *nameGV =
-      dyn_cast<GlobalVariable>(nameField->getOperand(0));
-    const ConstantDataArray *nameArray =
-      dyn_cast<ConstantDataArray>(nameGV->getInitializer());
-    // Lets add this string to the set of strings we should ignore processing
-    mIgnoreStr.insert(nameGV->getName());
-    if (mConstMems.find(nameGV->getName())
-        != mConstMems.end()) {
-      // If we already processesd this string as a constant, lets remove it from
-      // the list of known constants.  This way we don't process unneeded data
-      // and don't generate code/metadata for strings that are never used.
-      mConstMems.erase(mConstMems.find(nameGV->getName()));
-    } else {
-      mIgnoreStr.insert(CS->getOperand(0)->getName());
-    }
-    constPtr constAttr;
-    constAttr.name = nameArray->getAsString();
-    constAttr.size = (sizeField->getZExtValue() + 15) & ~15;
-    constAttr.base = CS;
-    constAttr.isArgument = true;
-    constAttr.isArray = false;
-    constAttr.cbNum = 0;
-    constAttr.offset = 0;
-    constAttr.usesHardware = (constAttr.size <= mSTM->device()->getMaxCBSize());
-    // Now that we have all our constant information,
-    // lets update the kernel
-    llvm::StringRef  kernelName = G->getName().data() + 30;
-    kernel k;
-    if (mKernels.find(kernelName) != mKernels.end()) {
-      k = mKernels[kernelName];
-    } else {
-      k.curSize = 0;
-      k.curRSize = 0;
-      k.curHWSize = 0;
-      k.curHWRSize = 0;
-      k.constSize = 0;
-      k.lvgv = NULL;
-      k.sgv = NULL;
-      memset(k.constSizes, 0, sizeof(uint32_t) * HW_MAX_NUM_CB);
-    }
-    constAttr.cbNum = k.constPtr.size() + 2;
-    k.constPtr.push_back(constAttr);
-    mKernels[kernelName] = k;
-  }
-}
-
-void AMDILGlobalManager::parseImageAnnotate(const GlobalValue *G) {
-  const GlobalVariable *GV = dyn_cast<GlobalVariable>(G);
-  const ConstantArray *CA = dyn_cast<ConstantArray>(GV->getInitializer());
-  if (!CA) {
-    return;
-  }
-  if (isa<GlobalValue>(CA)) {
-    return;
-  }
-  uint32_t e = CA->getNumOperands();
-  if (!e) {
-    return;
-  }
-  kernel k;
-  llvm::StringRef name = G->getName().data() + 23;
-  if (mKernels.find(name) != mKernels.end()) {
-    k = mKernels[name];
-  } else {
-    k.curSize = 0;
-    k.curRSize = 0;
-    k.curHWSize = 0;
-    k.curHWRSize = 0;
-    k.constSize = 0;
-    k.lvgv = NULL;
-    k.sgv = NULL;
-    memset(k.constSizes, 0, sizeof(uint32_t) * HW_MAX_NUM_CB);
-  }
-  for (uint32_t i = 0; i != e; ++i) {
-    const Value *V = CA->getOperand(i);
-    const Constant *C = dyn_cast<Constant>(V);
-    const ConstantStruct *CS = dyn_cast<ConstantStruct>(C);
-    if (CS && CS->getNumOperands() == 2) {
-      if (mConstMems.find(CS->getOperand(0)->getOperand(0)->getName()) !=
-          mConstMems.end()) {
-        // If we already processesd this string as a constant, lets remove it
-        // from the list of known constants.  This way we don't process unneeded
-        // data and don't generate code/metadata for strings that are never
-        // used.
-        mConstMems.erase(
-            mConstMems.find(CS->getOperand(0)->getOperand(0)->getName()));
-      } else {
-        mIgnoreStr.insert(CS->getOperand(0)->getOperand(0)->getName());
-      }
-      const ConstantInt *CI = dyn_cast<ConstantInt>(CS->getOperand(1));
-      uint32_t val = (uint32_t)CI->getZExtValue();
-      if (val == 1) {
-        k.readOnly.insert(i);
-      } else if (val == 2) {
-        k.writeOnly.insert(i);
-      } else {
-        assert(!"Unknown image type value!");
-      }
-    }
-  }
-  mKernels[name] = k;
-}
-
-void AMDILGlobalManager::parseAutoArray(const GlobalValue *GV, bool isRegion) {
-  const GlobalVariable *G = dyn_cast<GlobalVariable>(GV);
-  Type *Ty = (G) ? G->getType() : NULL;
-  arraymem tmp;
-  tmp.isHW = true;
-  tmp.offset = 0;
-  tmp.vecSize = getTypeSize(Ty, true);
-  tmp.isRegion = isRegion;
-  mArrayMems[GV->getName()] = tmp;
-}
-
-void AMDILGlobalManager::parseConstantPtr(const GlobalValue *GV) {
-  const GlobalVariable *G = dyn_cast<GlobalVariable>(GV);
-  Type *Ty = (G) ? G->getType() : NULL;
-  constPtr constAttr;
-  constAttr.name = G->getName();
-  constAttr.size = getTypeSize(Ty, true);
-  constAttr.base = GV;
-  constAttr.isArgument = false;
-  constAttr.isArray = true;
-  constAttr.offset = 0;
-  constAttr.cbNum = 0;
-  constAttr.usesHardware = false;
-  mConstMems[GV->getName()] = constAttr;
-}
-
-void AMDILGlobalManager::parseGlobalAnnotate(const GlobalValue *G) {
-  const GlobalVariable *GV = dyn_cast<GlobalVariable>(G);
-  if (!GV->hasInitializer()) {
-    return;
-  }
-  const Constant *CT = GV->getInitializer();
-  if (!CT || isa<GlobalValue>(CT)) {
-    return;
-  }
-  const ConstantArray *CA = dyn_cast<ConstantArray>(CT);
-  if (!CA) {
-    return;
-  }
-
-  unsigned int nKernels = CA->getNumOperands();
-  for (unsigned int i = 0, e = nKernels; i != e; ++i) {
-    parseKernelInformation(CA->getOperand(i));
-  }
-}
-
-void AMDILGlobalManager::parseKernelInformation(const Value *V) {
-  if (isa<GlobalValue>(V)) {
-    return;
-  }
-  const ConstantStruct *CS = dyn_cast_or_null<ConstantStruct>(V);
-  if (!CS) {
-    return;
-  }
-  uint32_t N = CS->getNumOperands();
-  if (N != 5) {
-    return;
-  }
-  kernel tmp;
-
-  tmp.curSize = 0;
-  tmp.curRSize = 0;
-  tmp.curHWSize = 0;
-  tmp.curHWRSize = 0;
-  // The first operand is always a pointer to the kernel.
-  const Constant *CV = dyn_cast<Constant>(CS->getOperand(0));
-  llvm::StringRef kernelName = "";
-  if (CV->getNumOperands()) {
-    kernelName = (*(CV->op_begin()))->getName();
-  }
-
-  // If we have images, then we have already created the kernel and we just need
-  // to get the kernel information.
-  if (mKernels.find(kernelName) != mKernels.end()) {
-    tmp = mKernels[kernelName];
-  } else {
-    tmp.curSize = 0;
-    tmp.curRSize = 0;
-    tmp.curHWSize = 0;
-    tmp.curHWRSize = 0;
-    tmp.constSize = 0;
-    tmp.lvgv = NULL;
-    tmp.sgv = NULL;
-    memset(tmp.constSizes, 0, sizeof(uint32_t) * HW_MAX_NUM_CB);
-  }
-
-
-  // The second operand is SGV, there can only be one so we don't need to worry
-  // about parsing out multiple data points.
-  CV = dyn_cast<Constant>(CS->getOperand(1));
-
-  llvm::StringRef sgvName;
-  if (CV->getNumOperands()) {
-    sgvName = (*(CV->op_begin()))->getName();
-  }
-
-  if (mKernelArgs.find(sgvName) != mKernelArgs.end()) {
-    tmp.sgv = &mKernelArgs[sgvName];
-  }
-  // The third operand is FGV, which is skipped
-  // The fourth operand is LVGV
-  // There can be multiple local arrays, so we
-  // need to handle each one seperatly
-  CV = dyn_cast<Constant>(CS->getOperand(3));
-  llvm::StringRef lvgvName = "";
-  if (CV->getNumOperands()) {
-    lvgvName = (*(CV->op_begin()))->getName();
-  }
-  if (mLocalArgs.find(lvgvName) != mLocalArgs.end()) {
-    localArg *ptr = &mLocalArgs[lvgvName];
-    tmp.lvgv = ptr;
-    llvm::SmallVector<arraymem *, DEFAULT_VEC_SLOTS>::iterator ib, ie;
-    for (ib = ptr->local.begin(), ie = ptr->local.end(); ib != ie; ++ib) {
-      if ((*ib)->isRegion) {
-        if ((*ib)->isHW) {
-          (*ib)->offset = tmp.curHWRSize;
-          tmp.curHWRSize += ((*ib)->vecSize + 15) & ~15;
-        } else {
-          (*ib)->offset = tmp.curRSize;
-          tmp.curRSize += ((*ib)->vecSize + 15) & ~15;
-        }
-      } else {
-        if ((*ib)->isHW) {
-          (*ib)->offset = tmp.curHWSize;
-          tmp.curHWSize += ((*ib)->vecSize + 15) & ~15;
-        } else {
-          (*ib)->offset = tmp.curSize;
-          tmp.curSize += ((*ib)->vecSize + 15) & ~15;
-        }
-      }
-    }
-  }
-
-  // The fifth operand is NULL
-  mKernels[kernelName] = tmp;
-}
-
-const kernel &AMDILGlobalManager::getKernel(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  assert(isKernel(name) && "Must be a kernel to call getKernel");
-  return iter->second;
-}
-
-bool AMDILGlobalManager::isKernel(const llvm::StringRef &name) const {
-  return (mKernels.find(name) != mKernels.end());
-}
-
-bool AMDILGlobalManager::isWriteOnlyImage(const llvm::StringRef &name,
-                                          uint32_t iID) const {
-  const StringMap<kernel>::const_iterator kiter = mKernels.find(name);
-  if (kiter == mKernels.end()) {
-    return false;
-  }
-  return kiter->second.writeOnly.count(iID);
-}
-
-uint32_t
-AMDILGlobalManager::getNumWriteImages(const llvm::StringRef &name) const {
-  char *env = NULL;
-  env = getenv("GPU_DISABLE_RAW_UAV");
-  if (env && env[0] == '1') {
-    return 8;
-  }
-  const StringMap<kernel>::const_iterator kiter = mKernels.find(name);
-  if (kiter == mKernels.end()) {
-    return 0;
-  } else {
-    return kiter->second.writeOnly.size();
-  }
-}
-
-bool AMDILGlobalManager::isReadOnlyImage(const llvm::StringRef &name,
-                                         uint32_t iID) const {
-  const StringMap<kernel>::const_iterator kiter = mKernels.find(name);
-  if (kiter == mKernels.end()) {
-    return false;
-  }
-  return kiter->second.readOnly.count(iID);
-}
-
-bool AMDILGlobalManager::hasRWG(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    kernelArg *ptr = iter->second.sgv;
-    if (ptr) {
-      return ptr->mHasRWG;
-    }
-  }
-  return false;
-}
-
-bool AMDILGlobalManager::hasRWR(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    kernelArg *ptr = iter->second.sgv;
-    if (ptr) {
-      return ptr->mHasRWR;
-    }
-  }
-  return false;
-}
-
-uint32_t
-AMDILGlobalManager::getMaxGroupSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    kernelArg *sgv = iter->second.sgv;
-    if (sgv) {
-      return sgv->reqGroupSize[0] * sgv->reqGroupSize[1] * sgv->reqGroupSize[2];
-    }
-  }
-  return mSTM->getDefaultSize(0) *
-         mSTM->getDefaultSize(1) *
-         mSTM->getDefaultSize(2);
-}
-
-uint32_t
-AMDILGlobalManager::getMaxRegionSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    kernelArg *sgv = iter->second.sgv;
-    if (sgv) {
-      return sgv->reqRegionSize[0] *
-             sgv->reqRegionSize[1] *
-             sgv->reqRegionSize[2];
-    }
-  }
-  return mSTM->getDefaultSize(0) *
-         mSTM->getDefaultSize(1) *
-         mSTM->getDefaultSize(2);
-}
-
-uint32_t AMDILGlobalManager::getRegionSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    return iter->second.curRSize;
-  } else {
-    return 0;
-  }
-}
-
-uint32_t AMDILGlobalManager::getLocalSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    return iter->second.curSize;
-  } else {
-    return 0;
-  }
-}
-
-uint32_t AMDILGlobalManager::getConstSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    return iter->second.constSize;
-  } else {
-    return 0;
-  }
-}
-
-uint32_t
-AMDILGlobalManager::getHWRegionSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    return iter->second.curHWRSize;
-  } else {
-    return 0;
-  }
-}
-
-uint32_t AMDILGlobalManager::getHWLocalSize(const llvm::StringRef &name) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end()) {
-    return iter->second.curHWSize;
-  } else {
-    return 0;
-  }
-}
-
-int32_t AMDILGlobalManager::getArgID(const Argument *arg) {
-  DenseMap<const Argument *, int32_t>::iterator argiter = mArgIDMap.find(arg);
-  if (argiter != mArgIDMap.end()) {
-    return argiter->second;
-  } else {
-    return -1;
-  }
-}
-
-
-uint32_t
-AMDILGlobalManager::getLocal(const llvm::StringRef &name, uint32_t dim) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end() && iter->second.sgv) {
-    kernelArg *sgv = iter->second.sgv;
-    switch (dim) {
-    default: break;
-    case 0:
-    case 1:
-    case 2:
-      return sgv->reqGroupSize[dim];
-      break;
-    case 3:
-      return sgv->reqGroupSize[0] * sgv->reqGroupSize[1] * sgv->reqGroupSize[2];
-    };
-  }
-  switch (dim) {
-  default:
-    return 1;
-  case 3:
-    return mSTM->getDefaultSize(0) *
-           mSTM->getDefaultSize(1) *
-           mSTM->getDefaultSize(2);
-  case 2:
-  case 1:
-  case 0:
-    return mSTM->getDefaultSize(dim);
-    break;
-  };
-  return 1;
-}
-
-uint32_t
-AMDILGlobalManager::getRegion(const llvm::StringRef &name, uint32_t dim) const {
-  StringMap<kernel>::const_iterator iter = mKernels.find(name);
-  if (iter != mKernels.end() && iter->second.sgv) {
-    kernelArg *sgv = iter->second.sgv;
-    switch (dim) {
-    default: break;
-    case 0:
-    case 1:
-    case 2:
-      return sgv->reqRegionSize[dim];
-      break;
-    case 3:
-      return sgv->reqRegionSize[0] *
-             sgv->reqRegionSize[1] *
-             sgv->reqRegionSize[2];
-    };
-  }
-  switch (dim) {
-  default:
-    return 1;
-  case 3:
-    return mSTM->getDefaultSize(0) *
-           mSTM->getDefaultSize(1) *
-           mSTM->getDefaultSize(2);
-  case 2:
-  case 1:
-  case 0:
-    return mSTM->getDefaultSize(dim);
-    break;
-  };
-  return 1;
-}
-
-StringMap<constPtr>::iterator AMDILGlobalManager::consts_begin() {
-  return mConstMems.begin();
-}
-
-
-StringMap<constPtr>::iterator AMDILGlobalManager::consts_end() {
-  return mConstMems.end();
-}
-
-bool AMDILGlobalManager::byteStoreExists(StringRef S) const {
-  return mByteStore.find(S) != mByteStore.end();
-}
-
-bool AMDILGlobalManager::usesHWConstant(const kernel &krnl,
-                                        const llvm::StringRef &arg) {
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->usesHardware;
-  } else {
-    return false;
-  }
-}
-
-uint32_t AMDILGlobalManager::getConstPtrSize(const kernel &krnl,
-                                             const llvm::StringRef &arg)
-{
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->size;
-  } else {
-    return 0;
-  }
-}
-
-uint32_t AMDILGlobalManager::getConstPtrOff(const kernel &krnl,
-                                            const llvm::StringRef &arg)
-{
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->offset;
-  } else {
-    return 0;
-  }
-}
-
-uint32_t AMDILGlobalManager::getConstPtrCB(const kernel &krnl,
-                                           const llvm::StringRef &arg)
-{
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->cbNum;
-  } else {
-    return 0;
-  }
-}
-
-void AMDILGlobalManager::calculateCPOffsets(const MachineFunction *MF,
-                                            kernel &krnl)
-{
-  const MachineConstantPool *MCP = MF->getConstantPool();
-  if (!MCP) {
-    return;
-  }
-  const std::vector<MachineConstantPoolEntry> consts = MCP->getConstants();
-  size_t numConsts = consts.size();
-  for (size_t x = 0; x < numConsts; ++x) {
-    krnl.CPOffsets.push_back(
-        std::make_pair<uint32_t, const Constant*>(
-          mCurrentCPOffset, consts[x].Val.ConstVal));
-    size_t curSize = getTypeSize(consts[x].Val.ConstVal->getType(), true);
-    // Align the size to the vector boundary
-    curSize = (curSize + 15) & (~15);
-    mCurrentCPOffset += curSize;
-  }
-}
-
-bool AMDILGlobalManager::isConstPtrArray(const kernel &krnl,
-                                         const llvm::StringRef &arg) {
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->isArray;
-  } else {
-    return false;
-  }
-}
-
-bool AMDILGlobalManager::isConstPtrArgument(const kernel &krnl,
-                                            const llvm::StringRef &arg)
-{
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->isArgument;
-  } else {
-    return false;
-  }
-}
-
-const Value *AMDILGlobalManager::getConstPtrValue(const kernel &krnl,
-                                                  const llvm::StringRef &arg) {
-  const constPtr *curConst = getConstPtr(krnl, arg);
-  if (curConst) {
-    return curConst->base;
-  } else {
-    return NULL;
-  }
-}
-
-static void
-dumpZeroElements(const  StructType * const T, llvm::raw_ostream &O, bool asBytes);
-static void
-dumpZeroElements(const IntegerType * const T, llvm::raw_ostream &O, bool asBytes);
-static void
-dumpZeroElements(const   ArrayType * const T, llvm::raw_ostream &O, bool asBytes);
-static void
-dumpZeroElements(const  VectorType * const T, llvm::raw_ostream &O, bool asBytes);
-static void
-dumpZeroElements(const        Type * const T, llvm::raw_ostream &O, bool asBytes);
-
-void dumpZeroElements(const Type * const T, llvm::raw_ostream &O, bool asBytes) {
-  if (!T) {
-    return;
-  }
-  switch(T->getTypeID()) {
-  case Type::X86_FP80TyID:
-  case Type::FP128TyID:
-  case Type::PPC_FP128TyID:
-  case Type::LabelTyID:
-    assert(0 && "These types are not supported by this backend");
-  default:
-  case Type::DoubleTyID:
-    if (asBytes) {
-      O << ":0:0:0:0:0:0:0:0";
-    } else {
-      O << ":0";
-    }
-    break;
-  case Type::FloatTyID:
-  case Type::PointerTyID:
-  case Type::FunctionTyID:
-    if (asBytes) {
-      O << ":0:0:0:0";
-    } else {
-      O << ":0";
-    }
-    break;
-  case Type::IntegerTyID:
-    dumpZeroElements(dyn_cast<IntegerType>(T), O, asBytes);
-    break;
-  case Type::StructTyID:
-    {
-      const StructType *ST = cast<StructType>(T);
-      if (!ST->isOpaque()) {
-        dumpZeroElements(dyn_cast<StructType>(T), O, asBytes);
-      } else { // A pre-LLVM 3.0 opaque type
-        if (asBytes) {
-          O << ":0:0:0:0";
-        } else {
-          O << ":0";
-        }
-      }
-    }
-    break;
-  case Type::ArrayTyID:
-    dumpZeroElements(dyn_cast<ArrayType>(T), O, asBytes);
-    break;
-  case Type::VectorTyID:
-    dumpZeroElements(dyn_cast<VectorType>(T), O, asBytes);
-    break;
-  };
-}
-
-void
-dumpZeroElements(const StructType * const ST, llvm::raw_ostream &O, bool asBytes) {
-  if (!ST) {
-    return;
-  }
-  Type *curType;
-  StructType::element_iterator eib = ST->element_begin();
-  StructType::element_iterator eie = ST->element_end();
-  for (;eib != eie; ++eib) {
-    curType = *eib;
-    dumpZeroElements(curType, O, asBytes);
-  }
-}
-
-void
-dumpZeroElements(const IntegerType * const IT, llvm::raw_ostream &O, bool asBytes) {
-  if (asBytes) {
-    unsigned byteWidth = (IT->getBitWidth() >> 3);
-    for (unsigned x = 0; x < byteWidth; ++x) {
-      O << ":0";
-    }
-  }
-}
-
-void
-dumpZeroElements(const ArrayType * const AT, llvm::raw_ostream &O, bool asBytes) {
-  size_t size = AT->getNumElements();
-  for (size_t x = 0; x < size; ++x) {
-    dumpZeroElements(AT->getElementType(), O, asBytes);
-  }
-}
-
-void
-dumpZeroElements(const VectorType * const VT, llvm::raw_ostream &O, bool asBytes) {
-  size_t size = VT->getNumElements();
-  for (size_t x = 0; x < size; ++x) {
-    dumpZeroElements(VT->getElementType(), O, asBytes);
-  }
-}
-
-void AMDILGlobalManager::printConstantValue(const Constant *CAval,
-                                            llvm::raw_ostream &O, bool asBytes) {
-  if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CAval)) {
-    bool isDouble = &CFP->getValueAPF().getSemantics()==&APFloat::IEEEdouble;
-    if (isDouble) {
-      double val = CFP->getValueAPF().convertToDouble();
-      union dtol_union {
-        double d;
-        uint64_t l;
-        char c[8];
-      } conv;
-      conv.d = val;
-      if (!asBytes) {
-        O << ":";
-        O.write_hex(conv.l);
-      } else {
-        for (int i = 0; i < 8; ++i) {
-          O << ":";
-          O.write_hex((unsigned)conv.c[i] & 0xFF);
-        }
-      }
-    } else {
-      float val = CFP->getValueAPF().convertToFloat();
-      union ftoi_union {
-        float f;
-        uint32_t u;
-        char c[4];
-      } conv;
-      conv.f = val;
-      if (!asBytes) {
-        O << ":";
-        O.write_hex(conv.u);
-      } else {
-        for (int i = 0; i < 4; ++i) {
-          O << ":";
-          O.write_hex((unsigned)conv.c[i] & 0xFF);
-        }
-      }
-    }
-  } else if (const ConstantInt *CI = dyn_cast<ConstantInt>(CAval)) {
-    uint64_t zVal = CI->getValue().getZExtValue();
-    if (!asBytes) {
-      O << ":";
-      O.write_hex(zVal);
-    } else {
-      switch (CI->getBitWidth()) {
-      default:
-        {
-          union ltob_union {
-            uint64_t l;
-            char c[8];
-          } conv;
-          conv.l = zVal;
-          for (int i = 0; i < 8; ++i) {
-            O << ":";
-            O.write_hex((unsigned)conv.c[i] & 0xFF);
-          }
-        }
-        break;
-      case 8:
-        O << ":";
-        O.write_hex(zVal & 0xFF);
-        break;
-      case 16:
-        {
-          union stob_union {
-            uint16_t s;
-            char c[2];
-          } conv;
-          conv.s = (uint16_t)zVal;
-          O << ":";
-          O.write_hex((unsigned)conv.c[0] & 0xFF);
-          O << ":";
-          O.write_hex((unsigned)conv.c[1] & 0xFF);
-        }
-        break;
-      case 32:
-        {
-          union itob_union {
-            uint32_t i;
-            char c[4];
-          } conv;
-          conv.i = (uint32_t)zVal;
-          for (int i = 0; i < 4; ++i) {
-            O << ":";
-            O.write_hex((unsigned)conv.c[i] & 0xFF);
-          }
-        }
-        break;
-      }
-    }
-  } else if (const ConstantVector *CV = dyn_cast<ConstantVector>(CAval)) {
-    int y = CV->getNumOperands()-1;
-    int x = 0;
-    for (; x < y; ++x) {
-      printConstantValue(CV->getOperand(x), O, asBytes);
-    }
-    printConstantValue(CV->getOperand(x), O, asBytes);
-  } else if (const ConstantStruct *CS = dyn_cast<ConstantStruct>(CAval)) {
-    int y = CS->getNumOperands();
-    int x = 0;
-    for (; x < y; ++x) {
-      printConstantValue(CS->getOperand(x), O, asBytes);
-    }
-  } else if (const ConstantAggregateZero *CAZ
-      = dyn_cast<ConstantAggregateZero>(CAval)) {
-    int y = CAZ->getNumOperands();
-    if (y > 0) {
-      int x = 0;
-      for (; x < y; ++x) {
-        printConstantValue((llvm::Constant *)CAZ->getOperand(x),
-            O, asBytes);
-      }
-    } else {
-      if (asBytes) {
-        dumpZeroElements(CAval->getType(), O, asBytes);
-      } else {
-        int y = getNumElements(CAval->getType())-1;
-        for (int x = 0; x < y; ++x) {
-          O << ":0";
-        }
-        O << ":0";
-      }
-    }
-  } else if (const ConstantArray *CA = dyn_cast<ConstantArray>(CAval)) {
-    int y = CA->getNumOperands();
-    int x = 0;
-    for (; x < y; ++x) {
-      printConstantValue(CA->getOperand(x), O, asBytes);
-    }
-  } else if (dyn_cast<ConstantPointerNull>(CAval)) {
-    O << ":0";
-    //assert(0 && "Hit condition which was not expected");
-  } else if (dyn_cast<ConstantExpr>(CAval)) {
-    O << ":0";
-    //assert(0 && "Hit condition which was not expected");
-  } else if (dyn_cast<UndefValue>(CAval)) {
-    O << ":0";
-    //assert(0 && "Hit condition which was not expected");
-  } else {
-    assert(0 && "Hit condition which was not expected");
-  }
-}
-
-static bool isStruct(Type * const T)
-{
-  if (!T) {
-    return false;
-  }
-  switch (T->getTypeID()) {
-  default:
-    return false;
-  case Type::PointerTyID:
-    return isStruct(T->getContainedType(0));
-  case Type::StructTyID:
-    return true;
-  case Type::ArrayTyID:
-  case Type::VectorTyID:
-    return isStruct(dyn_cast<SequentialType>(T)->getElementType());
-  };
-
-}
-
-void AMDILGlobalManager::dumpDataToCB(llvm::raw_ostream &O, AMDILKernelManager *km,
-                                      uint32_t id) {
-  uint32_t size = 0;
-  for (StringMap<constPtr>::iterator cmb = consts_begin(),
-      cme = consts_end(); cmb != cme; ++cmb) {
-    if (id == cmb->second.cbNum) {
-      size += (cmb->second.size + 15) & (~15);
-    }
-  }
-  if (id == 0) {
-    O << ";#DATASTART:" << (size + mCurrentCPOffset) << "\n";
-    if (mCurrentCPOffset) {
-      for (StringMap<kernel>::iterator kcpb = mKernels.begin(),
-          kcpe = mKernels.end(); kcpb != kcpe; ++kcpb) {
-        const kernel& k = kcpb->second;
-        size_t numConsts = k.CPOffsets.size();
-        for (size_t x = 0; x < numConsts; ++x) {
-          size_t offset = k.CPOffsets[x].first;
-          const Constant *C = k.CPOffsets[x].second;
-          Type *Ty = C->getType();
-          size_t size = (isStruct(Ty) ? getTypeSize(Ty, true)
-                                      : getNumElements(Ty));
-          O << ";#" << km->getTypeName(Ty, symTab) << ":";
-          O << offset << ":" << size ;
-          printConstantValue(C, O, isStruct(Ty));
-          O << "\n";
-        }
-      }
-    }
-  } else {
-    O << ";#DATASTART:" << id << ":" << size << "\n";
-  }
-
-  for (StringMap<constPtr>::iterator cmb = consts_begin(), cme = consts_end();
-       cmb != cme; ++cmb) {
-    if (cmb->second.cbNum != id) {
-      continue;
-    }
-    const GlobalVariable *G = dyn_cast<GlobalVariable>(cmb->second.base);
-    Type *Ty = (G) ? G->getType() : NULL;
-    size_t offset = cmb->second.offset;
-    const Constant *C = G->getInitializer();
-    size_t size = (isStruct(Ty)
-        ? getTypeSize(Ty, true)
-        : getNumElements(Ty));
-    O << ";#" << km->getTypeName(Ty, symTab) << ":";
-    if (!id) {
-      O << (offset + mCurrentCPOffset) << ":" << size;
-    } else {
-      O << offset << ":" << size;
-    }
-    if (C) {
-      printConstantValue(C, O, isStruct(Ty));
-    } else {
-      assert(0 && "Cannot have a constant pointer"
-          " without an initializer!");
-    }
-    O <<"\n";
-  }
-  if (id == 0) {
-    O << ";#DATAEND\n";
-  } else {
-    O << ";#DATAEND:" << id << "\n";
-  }
-}
-
-void
-AMDILGlobalManager::dumpDataSection(llvm::raw_ostream &O, AMDILKernelManager *km) {
-  if (mConstMems.empty() && !mCurrentCPOffset) {
-    return;
-  } else {
-    llvm::DenseSet<uint32_t> const_set;
-    for (StringMap<constPtr>::iterator cmb = consts_begin(), cme = consts_end();
-         cmb != cme; ++cmb) {
-      const_set.insert(cmb->second.cbNum);
-    }
-    if (mCurrentCPOffset) {
-      const_set.insert(0);
-    }
-    for (llvm::DenseSet<uint32_t>::iterator setb = const_set.begin(),
-           sete = const_set.end(); setb != sete; ++setb) {
-      dumpDataToCB(O, km, *setb);
-    }
-  }
-}
-
-/// Create a function ID if it is not known or return the known
-/// function ID.
-uint32_t AMDILGlobalManager::getOrCreateFunctionID(const GlobalValue* func) {
-  if (func->getName().size()) {
-    return getOrCreateFunctionID(func->getName());
-  } 
-  uint32_t id;
-  if (mFuncPtrNames.find(func) == mFuncPtrNames.end()) {
-    id = mFuncPtrNames.size() + RESERVED_FUNCS + mFuncNames.size();
-    mFuncPtrNames[func] = id;
-  } else {
-    id = mFuncPtrNames[func];
-  }
-  return id;
-}
-uint32_t AMDILGlobalManager::getOrCreateFunctionID(const std::string &func) {
-  uint32_t id;
-  if (mFuncNames.find(func) == mFuncNames.end()) {
-    id = mFuncNames.size() + RESERVED_FUNCS + mFuncPtrNames.size();
-    mFuncNames[func] = id;
-  } else {
-    id = mFuncNames[func];
-  }
-  return id;
-}
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_
diff --git a/src/gallium/drivers/radeon/AMDILISelLowering.cpp b/src/gallium/drivers/radeon/AMDILISelLowering.cpp
index 6f78d15ad0b..0f76babb807 100644
--- a/src/gallium/drivers/radeon/AMDILISelLowering.cpp
+++ b/src/gallium/drivers/radeon/AMDILISelLowering.cpp
@@ -14,9 +14,7 @@
 
 #include "AMDILISelLowering.h"
 #include "AMDILDevices.h"
-#include "AMDILGlobalManager.h"
 #include "AMDILIntrinsicInfo.h"
-#include "AMDILKernelManager.h"
 #include "AMDILMachineFunctionInfo.h"
 #include "AMDILSubtarget.h"
 #include "AMDILTargetMachine.h"
@@ -31,6 +29,7 @@
 #include "llvm/DerivedTypes.h"
 #include "llvm/Instructions.h"
 #include "llvm/Intrinsics.h"
+#include "llvm/Support/raw_ostream.h"
 #include "llvm/Target/TargetOptions.h"
 
 using namespace llvm;
@@ -1905,11 +1904,6 @@ AMDILTargetLowering::getTgtMemIntrinsic(IntrinsicInfo &Info,
              isRet = false;
              IntNo = AMDILISD::APPEND_CONSUME_NORET; break;
   };
-  const AMDILSubtarget *stm = &this->getTargetMachine()
-    .getSubtarget<AMDILSubtarget>();
-  AMDILKernelManager *KM = const_cast<AMDILKernelManager*>(
-      stm->getKernelManager());
-  KM->setOutputInst();
 
   Info.opc = IntNo;
   Info.memVT = (bitCastToInt) ? MVT::f32 : MVT::i32;
@@ -2134,58 +2128,33 @@ AMDILTargetLowering::LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const
   SDValue DST = Op;
   const GlobalAddressSDNode *GADN = cast<GlobalAddressSDNode>(Op);
   const GlobalValue *G = GADN->getGlobal();
-  const AMDILSubtarget *stm = &this->getTargetMachine()
-    .getSubtarget<AMDILSubtarget>();
-  const AMDILGlobalManager *GM = stm->getGlobalManager();
   DebugLoc DL = Op.getDebugLoc();
-  int64_t base_offset = GADN->getOffset();
-  int32_t arrayoffset = GM->getArrayOffset(G->getName());
-  int32_t constoffset = GM->getConstOffset(G->getName());
-  if (arrayoffset != -1) {
-    DST = DAG.getConstant(arrayoffset, MVT::i32);
-    DST = DAG.getNode(ISD::ADD, DL, MVT::i32,
-        DST, DAG.getConstant(base_offset, MVT::i32));
-  } else if (constoffset != -1) {
-    if (GM->getConstHWBit(G->getName())) {
-      DST = DAG.getConstant(constoffset, MVT::i32);
-      DST = DAG.getNode(ISD::ADD, DL, MVT::i32,
-          DST, DAG.getConstant(base_offset, MVT::i32));
-    } else {
-      SDValue addr = DAG.getTargetGlobalAddress(G, DL, MVT::i32);
-      SDValue DPReg = DAG.getRegister(AMDIL::SDP, MVT::i32);
-      DPReg = DAG.getNode(ISD::ADD, DL, MVT::i32, DPReg,
-          DAG.getConstant(base_offset, MVT::i32));
-      DST = DAG.getNode(AMDILISD::ADDADDR, DL, MVT::i32, addr, DPReg);
-    }
+  const GlobalVariable *GV = dyn_cast<GlobalVariable>(G);
+  if (!GV) {
+    DST = DAG.getTargetGlobalAddress(GV, DL, MVT::i32);
   } else {
-    const GlobalVariable *GV = dyn_cast<GlobalVariable>(G);
-    if (!GV) {
-      DST = DAG.getTargetGlobalAddress(GV, DL, MVT::i32);
-    } else {
-      if (GV->hasInitializer()) {
-        const Constant *C = dyn_cast<Constant>(GV->getInitializer());
-        if (const ConstantInt *CI = dyn_cast<ConstantInt>(C)) {
-          DST = DAG.getConstant(CI->getValue(), Op.getValueType());
-
-        } else if (const ConstantFP *CF = dyn_cast<ConstantFP>(C)) {
-          DST = DAG.getConstantFP(CF->getValueAPF(),
-              Op.getValueType());
-        } else if (dyn_cast<ConstantAggregateZero>(C)) {
-          EVT VT = Op.getValueType();
-          if (VT.isInteger()) {
-            DST = DAG.getConstant(0, VT);
-          } else {
-            DST = DAG.getConstantFP(0, VT);
-          }
+    if (GV->hasInitializer()) {
+      const Constant *C = dyn_cast<Constant>(GV->getInitializer());
+      if (const ConstantInt *CI = dyn_cast<ConstantInt>(C)) {
+        DST = DAG.getConstant(CI->getValue(), Op.getValueType());
+      } else if (const ConstantFP *CF = dyn_cast<ConstantFP>(C)) {
+        DST = DAG.getConstantFP(CF->getValueAPF(),
+            Op.getValueType());
+      } else if (dyn_cast<ConstantAggregateZero>(C)) {
+        EVT VT = Op.getValueType();
+        if (VT.isInteger()) {
+          DST = DAG.getConstant(0, VT);
         } else {
-          assert(!"lowering this type of Global Address "
-              "not implemented yet!");
-          C->dump();
-          DST = DAG.getTargetGlobalAddress(GV, DL, MVT::i32);
+          DST = DAG.getConstantFP(0, VT);
         }
       } else {
+        assert(!"lowering this type of Global Address "
+            "not implemented yet!");
+        C->dump();
         DST = DAG.getTargetGlobalAddress(GV, DL, MVT::i32);
       }
+    } else {
+      DST = DAG.getTargetGlobalAddress(GV, DL, MVT::i32);
     }
   }
   return DST;
diff --git a/src/gallium/drivers/radeon/AMDILKernelManager.cpp b/src/gallium/drivers/radeon/AMDILKernelManager.cpp
deleted file mode 100644
index 4df81ff5078..00000000000
--- a/src/gallium/drivers/radeon/AMDILKernelManager.cpp
+++ /dev/null
@@ -1,1356 +0,0 @@
-//===-- AMDILKernelManager.cpp - 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.
-//
-//==-----------------------------------------------------------------------===//
-#include "AMDILKernelManager.h"
-
-#include "AMDILAlgorithms.tpp"
-#include "AMDILKernelManager.h"
-#ifdef UPSTREAM_LLVM
-#include "AMDILAsmPrinter.h"
-#endif
-#include "AMDILCompilerErrors.h"
-#include "AMDILDeviceInfo.h"
-#include "AMDILDevices.h"
-#include "AMDILGlobalManager.h"
-#include "AMDILMachineFunctionInfo.h"
-#include "AMDILModuleInfo.h"
-#include "AMDILSubtarget.h"
-#include "AMDILTargetMachine.h"
-#include "AMDILUtilityFunctions.h"
-#include "llvm/ADT/StringExtras.h"
-#include "llvm/CodeGen/MachineFrameInfo.h"
-#include "llvm/Constants.h"
-#include "llvm/DerivedTypes.h"
-#include "llvm/Function.h"
-#include "llvm/Instructions.h"
-#include "llvm/Support/FormattedStream.h"
-#include "llvm/Support/MathExtras.h"
-
-#include <stdio.h>
-
-using namespace llvm;
-#define NUM_EXTRA_SLOTS_PER_IMAGE 1
-
-static bool errorPrint(const char *ptr, llvm::raw_ostream &O) {
-  if (ptr[0] == 'E') {
-    O << ";error:" << ptr << "\n";
-  } else {
-    O << ";warning:" << ptr << "\n";
-  }
-  return false;
-}
-
-#if 0
-static bool
-samplerPrint(StringMap<SamplerInfo>::iterator &data, llvm::raw_ostream &O) {
-  O << ";sampler:" << (*data).second.name << ":" << (*data).second.idx
-    << ":" << ((*data).second.val == (uint32_t)-1 ? 0 : 1) 
-    << ":" << ((*data).second.val != (uint32_t)-1 ? (*data).second.val : 0)
-    << "\n";
-  return false;
-}
-#endif
-
-static bool arenaPrint(uint32_t val, llvm::raw_ostream &O) {
-  if (val >= ARENA_SEGMENT_RESERVED_UAVS) {
-    O << "dcl_arena_uav_id(" << val << ")\n";
-  }
-  return false;
-}
-
-static bool uavPrint(uint32_t val, llvm::raw_ostream &O) {
-  if (val < 8 || val == 11){
-    O << "dcl_raw_uav_id(" << val << ")\n";
-  }
-  return false;
-}
-
-static bool uavPrintSI(uint32_t val, llvm::raw_ostream &O) {
-  O << "dcl_typeless_uav_id(" << val << ")_stride(4)_length(4)_access(read_write)\n";
-  return false;
-}
-
-static bool
-printfPrint(std::pair<const std::string, PrintfInfo *> &data, llvm::raw_ostream &O) {
-  O << ";printf_fmt:" << data.second->getPrintfID();
-  // Number of operands
-  O << ":" << data.second->getNumOperands();
-  // Size of each operand
-  for (size_t i = 0, e = data.second->getNumOperands(); i < e; ++i) {
-    O << ":" << (data.second->getOperandID(i) >> 3);
-  }
-  const char *ptr = data.first.c_str();
-  uint32_t size = data.first.size() - 1;
-  // The format string size
-  O << ":" << size << ":";
-  for (size_t i = 0; i < size; ++i) {
-    if (ptr[i] == '\r') {
-      O << "\\r";
-    } else if (ptr[i] == '\n') {
-      O << "\\n";
-    } else {
-      O << ptr[i];
-    }
-  }
-  O << ";\n";   // c_str() is cheap way to trim
-  return false;
-}
-
-
-void AMDILKernelManager::updatePtrArg(Function::const_arg_iterator Ip,
-                                      int numWriteImages, int raw_uav_buffer,
-                                      int counter, bool isKernel,
-                                      const Function *F) {
-  assert(F && "Cannot pass a NULL Pointer to F!");
-  assert(Ip->getType()->isPointerTy() &&
-         "Argument must be a pointer to be passed into this function!\n");
-  std::string ptrArg(";pointer:");
-  const char *symTab = "NoSymTab";
-  uint32_t ptrID = getUAVID(Ip);
-  const PointerType *PT = cast<PointerType>(Ip->getType());
-  uint32_t Align = 4;
-  const char *MemType = "uav";
-  if (PT->getElementType()->isSized()) {
-    Align = NextPowerOf2((uint32_t)mTM->getTargetData()->
-                            getTypeAllocSize(PT->getElementType()));
-  }
-  ptrArg += Ip->getName().str() + ":" + getTypeName(PT, symTab) + ":1:1:" +
-            itostr(counter * 16) + ":";
-  switch (PT->getAddressSpace()) {
-  case AMDILAS::ADDRESS_NONE:
-    //O << "No Address space qualifier!";
-    mMFI->addErrorMsg(amd::CompilerErrorMessage[INTERNAL_ERROR]);
-    assert(1);
-    break;
-  case AMDILAS::GLOBAL_ADDRESS:
-    if (mSTM->device()->isSupported(AMDILDeviceInfo::ArenaSegment)) {
-      if (ptrID >= ARENA_SEGMENT_RESERVED_UAVS) {
-        ptrID = 8;
-      }
-    }
-    mMFI->uav_insert(ptrID);
-    break;
-  case AMDILAS::CONSTANT_ADDRESS: {
-    if (isKernel && mSTM->device()->usesHardware(AMDILDeviceInfo::ConstantMem)){
-      const kernel t = mGM->getKernel(F->getName());
-      if (mGM->usesHWConstant(t, Ip->getName())) {
-        MemType = "hc\0";
-        ptrID = mGM->getConstPtrCB(t, Ip->getName());
-      } else {
-        MemType = "c\0";
-        mMFI->uav_insert(ptrID);
-      }
-    } else {
-      MemType = "c\0";
-      mMFI->uav_insert(ptrID);
-    }
-    break; 
-  }
-  default:
-  case AMDILAS::PRIVATE_ADDRESS:
-    if (mSTM->device()->usesHardware(AMDILDeviceInfo::PrivateMem)) {
-      MemType = (mSTM->device()->isSupported(AMDILDeviceInfo::PrivateUAV)) 
-        ? "up\0" : "hp\0";
-    } else {
-      MemType = "p\0";
-      mMFI->uav_insert(ptrID);
-    }
-    break;
-  case AMDILAS::REGION_ADDRESS:
-    mMFI->setUsesRegion();
-    if (mSTM->device()->usesHardware(AMDILDeviceInfo::RegionMem)) {
-      MemType = "hr\0";
-      ptrID = 0;
-    } else {
-      MemType = "r\0";
-      mMFI->uav_insert(ptrID);
-    }
-    break;
-  case AMDILAS::LOCAL_ADDRESS:
-    mMFI->setUsesLocal();
-    if (mSTM->device()->usesHardware(AMDILDeviceInfo::LocalMem)) {
-      MemType = "hl\0";
-      ptrID = 1;
-    } else {
-      MemType = "l\0";
-      mMFI->uav_insert(ptrID);
-    }
-    break;
-  };
-  ptrArg += std::string(MemType) + ":";
-  ptrArg += itostr(ptrID) + ":";
-  ptrArg += itostr(Align);
-  mMFI->addMetadata(ptrArg, true);
-}
-
-AMDILKernelManager::AMDILKernelManager(AMDILTargetMachine *TM,
-                                       AMDILGlobalManager *GM)
-{
-  mTM = TM;
-  mSTM = mTM->getSubtargetImpl();
-  mGM = GM;
-  clear();
-}
-
-AMDILKernelManager::~AMDILKernelManager() {
-  clear();
-}
-
-void 
-AMDILKernelManager::setMF(MachineFunction *MF)
-{
-  mMF = MF;
-  mMFI = MF->getInfo<AMDILMachineFunctionInfo>();
-}
-
-void AMDILKernelManager::clear() {
-  mUniqueID = 0;
-  mIsKernel = false;
-  mWasKernel = false;
-  mHasImageWrite = false;
-  mHasOutputInst = false;
-}
-
-bool AMDILKernelManager::useCompilerWrite(const MachineInstr *MI) {
-  return (MI->getOpcode() == AMDIL::RETURN && wasKernel() && !mHasImageWrite
-          && !mHasOutputInst);
-}
-
-void AMDILKernelManager::processArgMetadata(llvm::raw_ostream &O,
-                                            uint32_t buf,
-                                            bool isKernel) 
-{
-  const Function *F = mMF->getFunction();
-  const char * symTab = "NoSymTab";
-  Function::const_arg_iterator Ip = F->arg_begin();
-  Function::const_arg_iterator Ep = F->arg_end();
-  
-  if (F->hasStructRetAttr()) {
-    assert(Ip != Ep && "Invalid struct return fucntion!");
-    mMFI->addErrorMsg(amd::CompilerErrorMessage[INTERNAL_ERROR]);
-    ++Ip;
-  }
-  uint32_t mCBSize = 0;
-  int raw_uav_buffer = mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID);
-  bool MultiUAV = mSTM->device()->isSupported(AMDILDeviceInfo::MultiUAV);
-  bool ArenaSegment =
-    mSTM->device()->isSupported(AMDILDeviceInfo::ArenaSegment);
-  int numWriteImages =
-    mSTM->getGlobalManager()->getNumWriteImages(F->getName());
-  if (numWriteImages == OPENCL_MAX_WRITE_IMAGES || MultiUAV || ArenaSegment) {
-    if (mSTM->device()->getGeneration() <= AMDILDeviceInfo::HD6XXX) {
-      raw_uav_buffer = mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID);
-    }
-  }
-  uint32_t CounterNum = 0;
-  uint32_t ROArg = 0;
-  uint32_t WOArg = 0;
-  uint32_t NumArg = 0;
-  while (Ip != Ep) {
-    Type *cType = Ip->getType();
-    if (cType->isIntOrIntVectorTy() || cType->isFPOrFPVectorTy()) {
-      std::string argMeta(";value:");
-      argMeta += Ip->getName().str() + ":" + getTypeName(cType, symTab) + ":";
-      int bitsize = cType->getPrimitiveSizeInBits();
-      int numEle = 1;
-      if (cType->getTypeID() == Type::VectorTyID) {
-        numEle = cast<VectorType>(cType)->getNumElements();
-      }
-      argMeta += itostr(numEle) + ":1:" + itostr(mCBSize << 4);
-      mMFI->addMetadata(argMeta, true);
-
-      // FIXME: simplify
-      if ((bitsize / numEle) < 32) {
-        bitsize = numEle >> 2;
-      } else {
-        bitsize >>= 7;
-      }
-      if (!bitsize) {
-        bitsize = 1;
-      }
-
-      mCBSize += bitsize;
-      ++NumArg;
-    } else if (const PointerType *PT = dyn_cast<PointerType>(cType)) {
-      Type *CT = PT->getElementType();
-      const StructType *ST = dyn_cast<StructType>(CT);
-      if (ST && ST->isOpaque()) {
-        StringRef name = ST->getName();
-        bool i1d  = name.equals( "struct._image1d_t" );
-        bool i1da = name.equals( "struct._image1d_array_t" );
-        bool i1db = name.equals( "struct._image1d_buffer_t" );
-        bool i2d  = name.equals( "struct._image2d_t" );
-        bool i2da = name.equals( "struct._image2d_array_t" );
-        bool i3d  = name.equals( "struct._image3d_t" );
-        bool c32  = name.equals( "struct._counter32_t" );
-        bool c64  = name.equals( "struct._counter64_t" );
-        if (i1d || i1da || i1db || i2d | i2da || i3d) {
-          if (mSTM->device()->isSupported(AMDILDeviceInfo::Images)) {
-            std::string imageArg(";image:");
-            imageArg += Ip->getName().str() + ":";
-            if (i1d)       imageArg += "1D:";
-            else if (i1da) imageArg += "1DA:";
-            else if (i1db) imageArg += "1DB:";
-            else if (i2d)  imageArg += "2D:";
-            else if (i2da) imageArg += "2DA:";
-            else if (i3d)  imageArg += "3D:";
-
-            if (isKernel) {
-              if (mGM->isReadOnlyImage (mMF->getFunction()->getName(),
-                                        (ROArg + WOArg))) {
-                imageArg += "RO:" + itostr(ROArg);
-                O << "dcl_resource_id(" << ROArg << ")_type(";
-                if (i1d)       O << "1d";
-                else if (i1da) O << "1darray";
-                else if (i1db) O << "buffer";
-                else if (i2d)  O << "2d";
-                else if (i2da) O << "2darray";
-                else if (i3d)  O << "3d";
-                O << ")_fmtx(unknown)_fmty(unknown)"
-                  << "_fmtz(unknown)_fmtw(unknown)\n";
-                ++ROArg;
-              } else if (mGM->isWriteOnlyImage(mMF->getFunction()->getName(),
-                                               (ROArg + WOArg))) {
-                uint32_t offset = 0;
-                offset += WOArg;
-                imageArg += "WO:" + itostr(offset & 0x7);
-                O << "dcl_uav_id(" << ((offset) & 0x7) << ")_type(";
-                if (i1d)       O << "1d";
-                else if (i1da) O << "1darray";
-                else if (i1db) O << "buffer";
-                else if (i2d)  O << "2d";
-                else if (i2da) O << "2darray";
-                else if (i3d)  O << "3d";
-                O << ")_fmtx(uint)\n";
-                ++WOArg;
-              } else {
-                imageArg += "RW:" + itostr(ROArg + WOArg);
-              }
-            }
-            imageArg += ":1:" + itostr(mCBSize * 16);
-            mMFI->addMetadata(imageArg, true);
-            mMFI->addi32Literal(mCBSize);
-            mCBSize += NUM_EXTRA_SLOTS_PER_IMAGE + 1;
-            ++NumArg;
-          } else {
-            mMFI->addErrorMsg(amd::CompilerErrorMessage[NO_IMAGE_SUPPORT]);
-            ++NumArg;
-          }
-        } else if (c32 || c64) {
-          std::string counterArg(";counter:");
-          counterArg += Ip->getName().str() + ":"
-            + itostr(c32 ? 32 : 64) + ":"
-            + itostr(CounterNum++) + ":1:" + itostr(mCBSize * 16);
-          mMFI->addMetadata(counterArg, true);
-          ++NumArg;
-          ++mCBSize;
-        } else {
-          updatePtrArg(Ip, numWriteImages, raw_uav_buffer, mCBSize, isKernel,
-                       F);
-          ++NumArg;
-          ++mCBSize;
-        }
-      }
-        else if (CT->getTypeID() == Type::StructTyID
-                 && PT->getAddressSpace() == AMDILAS::PRIVATE_ADDRESS) {
-        const TargetData *td = mTM->getTargetData();
-        const StructLayout *sl = td->getStructLayout(dyn_cast<StructType>(CT));
-        int bytesize = sl->getSizeInBytes();
-        int reservedsize = (bytesize + 15) & ~15;
-        int numSlots = reservedsize >> 4;
-        if (!numSlots) {
-          numSlots = 1;
-        }
-        std::string structArg(";value:");
-        structArg += Ip->getName().str() + ":struct:"
-          + itostr(bytesize) + ":1:" + itostr(mCBSize * 16);
-        mMFI->addMetadata(structArg, true);
-        mCBSize += numSlots;
-        ++NumArg;
-      } else if (CT->isIntOrIntVectorTy()
-                 || CT->isFPOrFPVectorTy()
-                 || CT->getTypeID() == Type::ArrayTyID
-                 || CT->getTypeID() == Type::PointerTyID
-                 || PT->getAddressSpace() != AMDILAS::PRIVATE_ADDRESS) {
-        updatePtrArg(Ip, numWriteImages, raw_uav_buffer, mCBSize, isKernel, F);
-        ++NumArg;
-        ++mCBSize;
-      } else {
-        assert(0 && "Cannot process current pointer argument");
-        mMFI->addErrorMsg(amd::CompilerErrorMessage[INTERNAL_ERROR]);
-        ++NumArg;
-      }
-    } else {
-      assert(0 && "Cannot process current kernel argument");
-      mMFI->addErrorMsg(amd::CompilerErrorMessage[INTERNAL_ERROR]);
-      ++NumArg;
-    }
-    ++Ip;
-  }
-}
-
-void AMDILKernelManager::printHeader(AMDILAsmPrinter *AsmPrinter,
-                                     llvm::raw_ostream &O,
-                                     const std::string &name) {
-#ifdef UPSTREAM_LLVM
-  mName = name;
-  std::string kernelName;
-  kernelName = name;
-  int kernelId = mGM->getOrCreateFunctionID(kernelName);
-  O << "func " << kernelId << " ; " << kernelName << "\n";
-  if (mSTM->is64bit()) {
-    O << "mov " << AsmPrinter->getRegisterName(AMDIL::SDP) << ", cb0[8].xy\n";
-  } else {
-    O << "mov " << AsmPrinter->getRegisterName(AMDIL::SDP) << ", cb0[8].x\n";
-  }
-  O << "mov " << AsmPrinter->getRegisterName(AMDIL::SP) << ", l1.0\n";
-#endif
-}
-
-void AMDILKernelManager::printGroupSize(llvm::raw_ostream& O) {
-  // The HD4XXX generation of hardware does not support a 3D launch, so we need
-  // to use dcl_num_thread_per_group to specify the launch size. If the launch
-  // size is specified via a kernel attribute, we print it here. Otherwise we
-  // use the the default size.
-  if (mSTM->device()->getGeneration() == AMDILDeviceInfo::HD4XXX) {
-    if (mGM->hasRWG(mName) 
-        || !mMFI->usesLocal()) {
-      // if the user has specified what the required workgroup size is then we
-      // need to compile for that size and that size only.  Otherwise we compile
-      // for the max workgroup size that is passed in as an option to the
-      // backend.
-      O << "dcl_num_thread_per_group ";
-      O << mGM->getLocal(mName, 0) << ", ";
-      O << mGM->getLocal(mName, 1) << ", ";
-      O << mGM->getLocal(mName, 2) << "        \n";
-    } else {
-      // If the kernel uses local memory, then the kernel is being
-      // compiled in single wavefront mode. So we have to generate code slightly
-      // different.
-      O << "dcl_num_thread_per_group "
-        << mSTM->device()->getWavefrontSize()
-        << ", 1, 1       \n";
-    }
-  } else {
-    // Otherwise we generate for devices that support 3D launch natively.  If
-    // the reqd_workgroup_size attribute was specified, then we can specify the
-    // exact launch dimensions.
-    if (mGM->hasRWG(mName)) {
-      O << "dcl_num_thread_per_group ";
-      O << mGM->getLocal(mName, 0) << ", ";
-      O << mGM->getLocal(mName, 1) << ", ";
-      O << mGM->getLocal(mName, 2) << "        \n";
-    } else {
-      // Otherwise we specify the largest workgroup size that can be launched.
-      O << "dcl_max_thread_per_group " << mGM->getLocal(mName, 3) << " \n";
-    }
-  }
-  // Now that we have specified the workgroup size, lets declare the local
-  // memory size. If we are using hardware and we know the value at compile
-  // time, then we need to declare the correct value. Otherwise we should just
-  // declare the maximum size.
-  if (mSTM->device()->usesHardware(AMDILDeviceInfo::LocalMem)) {
-    size_t kernelLocalSize = (mGM->getHWLocalSize(mName) + 3) & ~3;
-    if (kernelLocalSize > mSTM->device()->getMaxLDSSize()) {
-      mMFI->addErrorMsg(amd::CompilerErrorMessage[INSUFFICIENT_LOCAL_RESOURCES]);
-    }
-    // If there is a local pointer as a kernel argument, we don't know the size
-    // at compile time, so we reserve all of the space.
-    if (mMFI->usesLocal() && (mMFI->hasLocalArg() || !kernelLocalSize)) {
-      O << "dcl_lds_id(" << DEFAULT_LDS_ID << ") "
-        << mSTM->device()->getMaxLDSSize() << "\n";
-      mMFI->setUsesMem(AMDILDevice::LDS_ID);
-    } else if (kernelLocalSize) {
-      // We know the size, so lets declare it correctly.
-      O << "dcl_lds_id(" << DEFAULT_LDS_ID << ") "
-        << kernelLocalSize << "\n";
-      mMFI->setUsesMem(AMDILDevice::LDS_ID);
-    }
-  }
-  // If the device supports the region memory extension, which maps to our
-  // hardware GDS memory, then lets declare it so we can use it later on.
-  if (mSTM->device()->usesHardware(AMDILDeviceInfo::RegionMem)) {
-    size_t kernelGDSSize = (mGM->getHWRegionSize(mName) + 3) & ~3;
-    if (kernelGDSSize > mSTM->device()->getMaxGDSSize()) {
-      mMFI->addErrorMsg(amd::CompilerErrorMessage[INSUFFICIENT_REGION_RESOURCES]);
-    }
-    // If there is a region pointer as a kernel argument, we don't know the size
-    // at compile time, so we reserved all of the space.
-    if (mMFI->usesRegion() && (mMFI->hasRegionArg() || !kernelGDSSize)) {
-      O << "dcl_gds_id(" << DEFAULT_GDS_ID <<
-        ") " << mSTM->device()->getMaxGDSSize() << "\n";
-      mMFI->setUsesMem(AMDILDevice::GDS_ID);
-    } else if (kernelGDSSize) {
-      // We know the size, so lets declare it.
-      O << "dcl_gds_id(" << DEFAULT_GDS_ID <<
-        ") " << kernelGDSSize << "\n";
-      mMFI->setUsesMem(AMDILDevice::GDS_ID);
-    }
-  }
-}
-
-void
-AMDILKernelManager::printDecls(AMDILAsmPrinter *AsmPrinter, llvm::raw_ostream &O) {
-  // If we are a HD4XXX generation device, then we only support a single uav
-  // surface, so we declare it and leave
-  if (mSTM->device()->getGeneration() == AMDILDeviceInfo::HD4XXX) {
-    O << "dcl_raw_uav_id(" 
-      << mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID)
-      << ")\n";
-    mMFI->setUsesMem(AMDILDevice::RAW_UAV_ID);
-    getIntrinsicSetup(AsmPrinter, O);
-    return;
-  }
-  // If we are supporting multiple uav's view the MultiUAV capability, then we
-  // need to print out the declarations here. MultiUAV conflicts with write
-  // images, so they only use 8 - NumWriteImages uav's. Therefor only pointers
-  // with ID's < 8 will get printed.
-  if (mSTM->device()->isSupported(AMDILDeviceInfo::MultiUAV)) {
-    binaryForEach(mMFI->uav_begin(), mMFI->uav_end(), uavPrint, O);
-    mMFI->setUsesMem(AMDILDevice::RAW_UAV_ID);
-  }
-  // If arena segments are supported, then we should emit them now.  Arena
-  // segments are similiar to MultiUAV, except ArenaSegments are virtual and up
-  // to 1024 of them can coexist. These are more compiler hints for CAL and thus
-  // cannot overlap in any form.  Each ID maps to a seperate piece of memory and
-  // CAL determines whether the load/stores should go to the fast path/slow path
-  // based on the usage and instruction.
-  if (mSTM->device()->isSupported(AMDILDeviceInfo::ArenaSegment)) {
-    binaryForEach(mMFI->uav_begin(), mMFI->uav_end(), arenaPrint, O);
-  }
-  // Now that we have printed out all of the arena and multi uav declaration,
-  // now we must print out the default raw uav id. This always exists on HD5XXX
-  // and HD6XXX hardware. The reason is that the hardware supports 12 UAV's and
-  // 11 are taken up by MultiUAV/Write Images and Arena.  However, if we do not
-  // have UAV 11 as the raw UAV and there are 8 write images, we must revert
-  // everything to the arena and not print out the default raw uav id.
-  if (mSTM->device()->getGeneration() == AMDILDeviceInfo::HD5XXX
-      || mSTM->device()->getGeneration() == AMDILDeviceInfo::HD6XXX) {
-    if ((mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID) < 11 &&
-         mSTM->getGlobalManager()->getNumWriteImages(mName)
-         != OPENCL_MAX_WRITE_IMAGES
-         && !mSTM->device()->isSupported(AMDILDeviceInfo::MultiUAV))
-        || mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID) == 11) {
-      if (!mMFI->usesMem(AMDILDevice::RAW_UAV_ID)
-          && mMFI->uav_count(mSTM->device()->
-              getResourceID(AMDILDevice::RAW_UAV_ID))) {
-        O << "dcl_raw_uav_id("
-          << mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID);
-        O << ")\n";
-        mMFI->setUsesMem(AMDILDevice::RAW_UAV_ID);
-      }
-    }
-    // If we have not printed out the arena ID yet, then do so here.
-      if (!mMFI->usesMem(AMDILDevice::ARENA_UAV_ID)
-          && mSTM->device()->usesHardware(AMDILDeviceInfo::ArenaUAV)) {
-        O << "dcl_arena_uav_id("
-          << mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID) << ")\n";
-        mMFI->setUsesMem(AMDILDevice::ARENA_UAV_ID);
-      }
-  } else if (mSTM->device()->getGeneration() > AMDILDeviceInfo::HD6XXX) {
-    binaryForEach(mMFI->uav_begin(), mMFI->uav_end(), uavPrintSI, O);
-    mMFI->setUsesMem(AMDILDevice::RAW_UAV_ID);
-  }
-  getIntrinsicSetup(AsmPrinter, O);
-}
-
-void AMDILKernelManager::getIntrinsicSetup(AMDILAsmPrinter *AsmPrinter,
-                                           llvm::raw_ostream &O)
-{
-  O << "mov r0.z, vThreadGrpIdFlat.x\n"
-    << "mov r1022.xyz0, vTidInGrp.xyz\n";
-  if (mSTM->device()->getGeneration() > AMDILDeviceInfo::HD4XXX) {
-    O << "mov r1023.xyz0, vThreadGrpId.xyz\n";
-   } else {
-    O << "imul r0.w, cb0[2].x, cb0[2].y\n"
-      // Calculates the local id.
-      // Calculates the group id.
-      << "umod r1023.x, r0.z, cb0[2].x\n"
-      << "udiv r1023.y, r0.z, cb0[2].x\n"
-      << "umod r1023.y, r1023.y, cb0[2].y\n"
-      << "udiv r1023.z, r0.z, r0.w\n";
-  }
-  // Calculates the global id.
-  if (mGM->hasRWG(mName) && 0) {
-    // Anytime we declare a literal, we need to reserve it, if it is not emitted
-    // in emitLiterals.
-    mMFI->addReservedLiterals(1);
-    O << "dcl_literal l" << mMFI->getNumLiterals() + 1 << ", ";
-    O << mGM->getLocal(mName, 0) << ", ";
-    O << mGM->getLocal(mName, 1) << ", ";
-    O << mGM->getLocal(mName, 2) << ", ";
-    O << "0\n";
-    O << "imad r1021.xyz0, r1023.xyz, l" << mMFI->getNumLiterals() + 1 << ".xyz, r1022.xyz\n";
-    mMFI->addReservedLiterals(1);
-  } else {
-    O << "imad r1021.xyz0, r1023.xyz, cb0[1].xyz, r1022.xyz\n";
-  }
-
-  // Add the global/group offset for multi-launch support.
-  O << "iadd r1021.xyz0, r1021.xyz0, cb0[6].xyz0\n"
-    << "iadd r1023.xyz0, r1023.xyz0, cb0[7].xyz0\n"
-    // moves the flat group id.
-    << "mov r1023.w, r0.z\n";
-#ifdef UPSTREAM_LLVM
-  if (mSTM->device()->usesSoftware(AMDILDeviceInfo::LocalMem)) {
-    if (mSTM->is64bit()) {
-      O << "umul " << AsmPrinter->getRegisterName(AMDIL::T2) 
-        << ".x0, r1023.w, cb0[4].z\n"
-        << "i64add " << AsmPrinter->getRegisterName(AMDIL::T2)
-        << ".xy, " << AsmPrinter->getRegisterName(AMDIL::T2)
-        << ".xy, cb0[4].xy\n";
-
-    } else {
-      O << "imad " << AsmPrinter->getRegisterName(AMDIL::T2)
-        << ".x, r1023.w, cb0[4].y, cb0[4].x\n";
-    }
-  }
-  // Shift the flat group id to be in bytes instead of dwords.
-  O << "ishl r1023.w, r1023.w, l0.z\n";
-  if (mSTM->device()->usesSoftware(AMDILDeviceInfo::PrivateMem)) {
-    if (mSTM->is64bit()) {
-      O << "umul " << AsmPrinter->getRegisterName(AMDIL::T1) 
-        << ".x0, vAbsTidFlat.x, cb0[3].z\n"
-        << "i64add " << AsmPrinter->getRegisterName(AMDIL::T1)
-        << ".xy, " << AsmPrinter->getRegisterName(AMDIL::T1)
-        << ".xy, cb0[3].xy\n";
-
-    } else {
-      O << "imad " << AsmPrinter->getRegisterName(AMDIL::T1)
-        << ".x, vAbsTidFlat.x, cb0[3].y, cb0[3].x\n";
-    }
-  } else {
-    O << "mov " << AsmPrinter->getRegisterName(AMDIL::T1) << ".x, l0.0\n";
-  }
-#endif
-  if (mSTM->device()->isSupported(AMDILDeviceInfo::RegionMem)) {
-    O << "udiv r1024.xyz, r1021.xyz, cb0[10].xyz\n";
-    if (mGM->hasRWR(mName) && 0) {
-      // Anytime we declare a literal, we need to reserve it, if it is not emitted
-      // in emitLiterals.
-      mMFI->addReservedLiterals(1);
-      O << "dcl_literal l" << mMFI->getNumLiterals() + 1 << ", ";
-      O << mGM->getLocal(mName, 0) << ", ";
-      O << mGM->getLocal(mName, 1) << ", ";
-      O << mGM->getLocal(mName, 2) << ", ";
-      O << "0\n";
-      O << "imad r1025.xyz0, r1023.xyz, l" << mMFI->getNumLiterals() + 1 << ".xyz, r1022.xyz\n";
-      mMFI->addReservedLiterals(1);
-    } else {
-      O << "imad r1025.xyz0, r1023.xyz, cb0[1].xyz, r1022.xyz\n";
-    }
-  }
-}
-
-void AMDILKernelManager::printFooter(llvm::raw_ostream &O) {
-  O << "ret\n";
-  O << "endfunc ; " << mName << "\n";
-}
-
-void
-AMDILKernelManager::printMetaData(llvm::raw_ostream &O, uint32_t id, bool kernel) {
-  if (kernel) {
-    int kernelId = mGM->getOrCreateFunctionID(mName);
-    mMFI->addCalledFunc(id);
-    mUniqueID = kernelId;
-    mIsKernel = true;
-  }
-  printKernelArgs(O);
-  if (kernel) {
-    mIsKernel = false;
-    mMFI->eraseCalledFunc(id);
-    mUniqueID = id;
-  }
-}
-
-void AMDILKernelManager::setKernel(bool kernel) {
-  mIsKernel = kernel;
-  if (kernel) {
-    mWasKernel = mIsKernel;
-  }
-}
-
-void AMDILKernelManager::setID(uint32_t id)
-{
-  mUniqueID = id;
-}
-
-void AMDILKernelManager::setName(const std::string &name) {
-  mName = name;
-}
-
-bool AMDILKernelManager::isKernel() {
-  return mIsKernel;
-}
-
-bool AMDILKernelManager::wasKernel() {
-  return mWasKernel;
-}
-
-void AMDILKernelManager::setImageWrite() {
-  mHasImageWrite = true;
-}
-
-void AMDILKernelManager::setOutputInst() {
-  mHasOutputInst = true;
-}
-
-void AMDILKernelManager::printConstantToRegMapping(
-       AMDILAsmPrinter *RegNames,
-       uint32_t &LII,
-       llvm::raw_ostream &O,
-       uint32_t &Counter,
-       uint32_t Buffer,
-       uint32_t n,
-       const char *lit,
-       uint32_t fcall,
-       bool isImage,
-       bool isHWCB)
-{
-#ifdef UPSTREAM_LLVM
-  // TODO: This needs to be enabled or SC will never statically index into the
-  // CB when a pointer is used.
-  if (mSTM->device()->usesHardware(AMDILDeviceInfo::ConstantMem) && isHWCB) {
-    const char *name = RegNames->getRegisterName(LII);
-    O << "mov " << name << ", l5.x\n";
-    ++LII;
-    Counter++;
-    return;
-  }
-  for (uint32_t x = 0; x < n; ++x) {
-    const char *name = RegNames->getRegisterName(LII);
-    if (isImage) {
-      O << "mov " << name << ", l" << mMFI->getIntLits(Counter++) << "\n";
-    } else {
-      O << "mov " << name << ", cb" <<Buffer<< "[" <<Counter++<< "]\n";
-    }
-    switch(fcall) {
-    case 1093:
-      O << "ishr " << name << ", " << name << ".xxyy, l3.0y0y\n"
-        "ishl " << name << ", " << name << ", l3.y\n"
-        "ishr " << name << ", " << name << ", l3.y\n";
-      break;
-    case 1092:
-      O << "ishr " << name << ", " << name << ".xx, l3.0y\n"
-        "ishl " << name << ", " << name << ", l3.y\n"
-        "ishr " << name << ", " << name << ", l3.y\n";
-      break;
-    case 1091:
-      O << "ishr " << name << ", " << name << ".xxxx, l3.0zyx\n"
-        "ishl " << name << ", " << name << ", l3.x\n"
-        "ishr " << name << ", " << name << ", l3.x\n";
-      break;
-    case 1090:
-      O << "ishr " << name << ", " << name << ".xx, l3.0z\n"
-        "ishl " << name << ".xy__, " << name << ".xy, l3.x\n"
-        "ishr " << name << ".xy__, " << name << ".xy, l3.x\n";
-      break;
-    default:
-      break;
-    };
-    if (lit) {
-      O << "ishl " << name << ", " << name
-        << ", " << lit << "\n";
-      O << "ishr " << name << ", " << name
-        << ", " << lit << "\n";
-    }
-    if (isImage) {
-      Counter += NUM_EXTRA_SLOTS_PER_IMAGE;
-    }
-    ++LII;
-  }
-#endif
-}
-
-void
-AMDILKernelManager::printCopyStructPrivate(const StructType *ST,
-                                           llvm::raw_ostream &O,
-                                           size_t stackSize,
-                                           uint32_t Buffer,
-                                           uint32_t mLitIdx,
-                                           uint32_t &Counter)
-{
-  size_t n = ((stackSize + 15) & ~15) >> 4;
-  for (size_t x = 0; x < n; ++x) {
-    O << "mov r2, cb" << Buffer << "[" << Counter++ << "]\n";
-    O << "mov r1.x, r0.x\n";
-    if (mSTM->device()->getGeneration() <= AMDILDeviceInfo::HD6XXX) {
-    if (mSTM->device()->usesHardware(AMDILDeviceInfo::PrivateMem)) {
-      O << "ishr r1.x, r1.x, l0.x\n";
-      O << "mov x" << mSTM->device()->getResourceID(AMDILDevice::SCRATCH_ID)
-        <<"[r1.x], r2\n";
-    } else {
-        O << "uav_raw_store_id(" <<
-          mSTM->device()->getResourceID(AMDILDevice::GLOBAL_ID)
-          << ") mem0, r1.x, r2\n";
-    }
-    } else {
-      O << "uav_raw_store_id(" <<
-        mSTM->device()->getResourceID(AMDILDevice::SCRATCH_ID)
-        << ") mem0, r1.x, r2\n";
-    }
-    O << "iadd r0.x, r0.x, l" << mLitIdx << ".z\n";
-  }
-}
-
-void AMDILKernelManager::printKernelArgs(llvm::raw_ostream &O) {
-  std::string version(";version:");
-  version += itostr(AMDIL_MAJOR_VERSION) + ":"
-    + itostr(AMDIL_MINOR_VERSION) + ":" + itostr(AMDIL_REVISION_NUMBER);
-  O << ";ARGSTART:" <<mName<< "\n";
-  if (mIsKernel) {
-    O << version << "\n";
-    O << ";device:" <<mSTM->getDeviceName() << "\n";
-  }
-  O << ";uniqueid:" <<mUniqueID<< "\n";
-  
-  size_t local = mGM->getLocalSize(mName);
-  size_t hwlocal = ((mGM->getHWLocalSize(mName) + 3) & (~0x3));
-  size_t region = mGM->getRegionSize(mName);
-  size_t hwregion = ((mGM->getHWRegionSize(mName) + 3) & (~0x3));
-  bool usehwlocal = mSTM->device()->usesHardware(AMDILDeviceInfo::LocalMem);
-  bool usehwprivate = mSTM->device()->usesHardware(AMDILDeviceInfo::PrivateMem);
-  bool usehwregion = mSTM->device()->usesHardware(AMDILDeviceInfo::RegionMem);
-  bool useuavprivate = mSTM->device()->isSupported(AMDILDeviceInfo::PrivateUAV);
-  if (mIsKernel) {
-    O << ";memory:" << ((usehwprivate) ? 
-        (useuavprivate) ? "uav" : "hw" : "" ) << "private:"
-      <<(((mMFI->getStackSize() + 15) & (~0xF)))<< "\n";
-  }
-  if (mSTM->device()->isSupported(AMDILDeviceInfo::RegionMem)) {
-    O << ";memory:" << ((usehwregion) ? "hw" : "") << "region:"
-      << ((usehwregion) ? hwregion : hwregion + region) << "\n";
-  }
-  O << ";memory:" << ((usehwlocal) ? "hw" : "") << "local:"
-    << ((usehwlocal) ? hwlocal : hwlocal + local) << "\n";
-  
-  if (mIsKernel) {
-    if (mGM->hasRWG(mName)) {
-      O << ";cws:" << mGM->getLocal(mName, 0) << ":";
-      O << mGM->getLocal(mName, 1) << ":";
-      O << mGM->getLocal(mName, 2) << "\n";
-    }
-    if (mGM->hasRWR(mName)) {
-      O << ";crs:" << mGM->getRegion(mName, 0) << ":";
-      O << mGM->getRegion(mName, 1) << ":";
-      O << mGM->getRegion(mName, 2) << "\n";
-    }
-  }
-  if (mIsKernel) {
-    for (std::vector<std::string>::iterator ib = mMFI->kernel_md_begin(),
-           ie = mMFI->kernel_md_end(); ib != ie; ++ib) {
-      O << (*ib) << "\n";
-    }
-  }
-  for (std::set<std::string>::iterator ib = mMFI->func_md_begin(),
-         ie = mMFI->func_md_end(); ib != ie; ++ib) {
-    O << (*ib) << "\n";
-  }
-  if (!mMFI->func_empty()) {
-    O << ";function:" << mMFI->func_size();
-    binaryForEach(mMFI->func_begin(), mMFI->func_end(), commaPrint, O);
-    O << "\n";
-  }
-
-  if (!mSTM->device()->isSupported(AMDILDeviceInfo::MacroDB)
-      && !mMFI->intr_empty()) {
-    O << ";intrinsic:" << mMFI->intr_size();
-    binaryForEach(mMFI->intr_begin(), mMFI->intr_end(), commaPrint, O);
-    O << "\n";
-  }
-
-  if (!mIsKernel) {
-    binaryForEach(mMFI->printf_begin(), mMFI->printf_end(), printfPrint, O);
-    mMF->getMMI().getObjFileInfo<AMDILModuleInfo>().add_printf_offset(
-        mMFI->printf_size());
-  } else {
-    for (StringMap<SamplerInfo>::iterator 
-        smb = mMFI->sampler_begin(),
-        sme = mMFI->sampler_end(); smb != sme; ++ smb) {
-      O << ";sampler:" << (*smb).second.name << ":" << (*smb).second.idx
-        << ":" << ((*smb).second.val == (uint32_t)-1 ? 0 : 1) 
-        << ":" << ((*smb).second.val != (uint32_t)-1 ? (*smb).second.val : 0)
-        << "\n";
-    }
-  }
-  if (mSTM->is64bit()) {
-    O << ";memory:64bitABI\n";
-  }
-
-  if (mMFI->errors_empty()) {
-    binaryForEach(mMFI->errors_begin(), mMFI->errors_end(), errorPrint, O);
-  }
-  // This has to come last
-  if (mIsKernel 
-      && mSTM->device()->getGeneration() <= AMDILDeviceInfo::HD6XXX) {
-    if (mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID) >
-        mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID)) {
-      if (mMFI->uav_size() == 1) {
-        if (mSTM->device()->isSupported(AMDILDeviceInfo::ArenaSegment)
-            && *(mMFI->uav_begin()) >= ARENA_SEGMENT_RESERVED_UAVS) {
-          O << ";uavid:"
-            << mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID);
-          O << "\n";
-        } else {
-          O << ";uavid:" << *(mMFI->uav_begin()) << "\n";
-        }
-      } else if (mMFI->uav_count(mSTM->device()->
-            getResourceID(AMDILDevice::RAW_UAV_ID))) {
-        O << ";uavid:"
-          << mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID);
-        O << "\n";
-      } else {
-        O << ";uavid:"
-          << mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID);
-        O << "\n";
-      }
-    } else if (mSTM->getGlobalManager()->getNumWriteImages(mName) !=
-        OPENCL_MAX_WRITE_IMAGES
-        && !mSTM->device()->isSupported(AMDILDeviceInfo::ArenaSegment)
-        && mMFI->uav_count(mSTM->device()->
-          getResourceID(AMDILDevice::RAW_UAV_ID))) {
-      O << ";uavid:"
-        << mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID) << "\n";
-    } else if (mMFI->uav_size() == 1) {
-      O << ";uavid:" << *(mMFI->uav_begin()) << "\n";
-    } else {
-      O << ";uavid:"
-        << mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID);
-      O << "\n";
-    }
-  }
-  O << ";ARGEND:" << mName << "\n";
-}
-
-void AMDILKernelManager::printArgCopies(llvm::raw_ostream &O,
-    AMDILAsmPrinter *RegNames)
-{
-  Function::const_arg_iterator I = mMF->getFunction()->arg_begin();
-  Function::const_arg_iterator Ie = mMF->getFunction()->arg_end();
-  uint32_t Counter = 0;
-
-  if (mMFI->getArgSize()) {
-    O << "dcl_cb cb1";
-    O << "[" << (mMFI->getArgSize() >> 4) << "]\n";
-    mMFI->setUsesMem(AMDILDevice::CONSTANT_ID);
-  }
-  const Function *F = mMF->getFunction();
-  // Get the stack size
-  uint32_t stackSize = mMFI->getStackSize();
-  uint32_t privateSize = mMFI->getScratchSize();
-  uint32_t stackOffset = (privateSize + 15) & (~0xF);
-  if (stackSize 
-      && mSTM->device()->usesHardware(AMDILDeviceInfo::PrivateMem)) {
-    // TODO: If the size is too large, we need to fall back to software emulated
-    // instead of using the hardware capability.
-    int size = (((stackSize + 15) & (~0xF)) >> 4);
-    if (size > 4096) {
-      mMFI->addErrorMsg(amd::CompilerErrorMessage[INSUFFICIENT_PRIVATE_RESOURCES]);
-    }
-    if (size) {
-    // For any stack variables, we need to declare the literals for them so that
-    // we can use them when we copy our data to the stack.
-    mMFI->addReservedLiterals(1);
-    // Anytime we declare a literal, we need to reserve it, if it is not emitted
-    // in emitLiterals.
-#ifdef UPSTREAM_LLVM
-    O << "dcl_literal l" << mMFI->getNumLiterals() << ", " << stackSize << ", "
-      << privateSize << ", 16, " << ((stackSize == privateSize) ? 0 : stackOffset) << "\n"
-      << "iadd r0.x, " << RegNames->getRegisterName(AMDIL::T1) << ".x, l"
-      << mMFI->getNumLiterals() << ".w\n";
-    if (mSTM->device()->getGeneration() <= AMDILDeviceInfo::HD6XXX) {
-    O << "dcl_indexed_temp_array x"
-      << mSTM->device()->getResourceID(AMDILDevice::SCRATCH_ID) << "["
-      << size << "]\n";
-    } else {
-      O << "dcl_typeless_uav_id("
-        << mSTM->device()->getResourceID(AMDILDevice::SCRATCH_ID) 
-        << ")_stride(4)_length(" << (size  << 4 )<< ")_access(private)\n";
-
-    }
-    O << "mov " << RegNames->getRegisterName(AMDIL::FP) 
-      << ".x, l" << mMFI->getNumLiterals() << ".0\n";
-#endif    
-    mMFI->setUsesMem(AMDILDevice::SCRATCH_ID);
-    }
-  }
-  I = mMF->getFunction()->arg_begin();
-  int32_t count = 0;
-  // uint32_t Image = 0;
-  bool displaced1 = false;
-  bool displaced2 = false;
-  uint32_t curReg = AMDIL::R1;
-  // TODO: We don't handle arguments that were pushed onto the stack!
-  for (; I != Ie; ++I) {
-    Type *curType = I->getType();
-    unsigned int Buffer = 1;
-    O << "; Kernel arg setup: " << I->getName() << "\n";
-    if (curType->isIntegerTy() || curType->isFloatingPointTy()) {
-      switch (curType->getPrimitiveSizeInBits()) {
-        default:
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1);
-          break;
-        case 16:
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1,
-              "l3.y" );
-          break;
-        case 8:
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1, "l3.x" );
-          break;
-      }
-#ifdef UPSTREAM_LLVM
-    } else if (const VectorType *VT = dyn_cast<VectorType>(curType)) {
-      Type *ET = VT->getElementType();
-      int numEle = VT->getNumElements();
-      switch (ET->getPrimitiveSizeInBits()) {
-        default:
-          if (numEle == 3) {
-            O << "mov " << RegNames->getRegisterName(curReg);
-            O << ".x, cb" << Buffer << "[" << Counter << "].x\n";
-            curReg++;
-            O << "mov " << RegNames->getRegisterName(curReg);
-            O << ".x, cb" << Buffer << "[" << Counter << "].y\n";
-            curReg++;
-            O << "mov " << RegNames->getRegisterName(curReg);
-            O << ".x, cb" << Buffer << "[" << Counter << "].z\n";
-            curReg++;
-            Counter++;
-          } else {
-            printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer,
-                (numEle+2) >> 2);
-          }
-          break;
-        case 64:
-          if (numEle == 3) {
-            O << "mov " << RegNames->getRegisterName(curReg);
-            O << ".xy, cb" << Buffer << "[" << Counter << "].xy\n";
-            curReg++;
-            O << "mov " << RegNames->getRegisterName(curReg);
-            O << ".xy, cb" << Buffer << "[" << Counter++ << "].zw\n";
-            curReg++;
-            O << "mov " << RegNames->getRegisterName(curReg);
-            O << ".xy, cb" << Buffer << "[" << Counter << "].xy\n";
-            curReg++;
-            Counter++;
-          } else {
-            printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer,
-                (numEle) >> 1);
-          }
-          break;
-        case 16: 
-          {
-                   switch (numEle) {
-                     default:
-                       printConstantToRegMapping(RegNames, curReg, O, Counter,
-                           Buffer, (numEle+2) >> 2, "l3.y", 1093);
-                       if (numEle == 3) {
-                         O << "mov " << RegNames->getRegisterName(curReg) << ".x, ";
-                         O << RegNames->getRegisterName(curReg) << ".y\n";
-                         ++curReg;
-                         O << "mov " << RegNames->getRegisterName(curReg) << ".x, ";
-                         O << RegNames->getRegisterName(curReg) << ".z\n";
-                         ++curReg;
-                       }
-                       break;
-                     case 2:
-                       printConstantToRegMapping(RegNames, curReg, O, Counter,
-                           Buffer, 1, "l3.y", 1092);
-                       break;
-                   }
-                   break;
-                 }
-        case 8: 
-          {
-                  switch (numEle) {
-                    default:
-                      printConstantToRegMapping(RegNames, curReg, O, Counter,
-                          Buffer, (numEle+2) >> 2, "l3.x", 1091);
-                      if (numEle == 3) {
-                        O << "mov " << RegNames->getRegisterName(curReg) << ".x, ";
-                        O << RegNames->getRegisterName(curReg) << ".y\n";
-                        ++curReg;
-                        O << "mov " << RegNames->getRegisterName(curReg) << ".x, ";
-                        O << RegNames->getRegisterName(curReg) << ".z\n";
-                        ++curReg;
-                      }
-                      break;
-                    case 2:
-                      printConstantToRegMapping(RegNames, curReg, O, Counter,
-                          Buffer, 1, "l3.x", 1090);
-                      break;
-                  }
-                  break;
-                }
-      }
-#endif
-    } else if (const PointerType *PT = dyn_cast<PointerType>(curType)) {
-      Type *CT = PT->getElementType();
-      const StructType *ST = dyn_cast<StructType>(CT);
-      if (ST && ST->isOpaque()) {
-        bool i1d  = ST->getName() == "struct._image1d_t";
-        bool i1da = ST->getName() == "struct._image1d_array_t";
-        bool i1db = ST->getName() == "struct._image1d_buffer_t";
-        bool i2d  = ST->getName() == "struct._image2d_t";
-        bool i2da = ST->getName() == "struct._image2d_array_t";
-        bool i3d  = ST->getName() == "struct._image3d_t";
-        bool is_image = i1d || i1da || i1db || i2d || i2da || i3d;
-        if (is_image) {
-          if (mSTM->device()->isSupported(AMDILDeviceInfo::Images)) {
-            printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer,
-                1, NULL, 0, is_image);
-          } else {
-            mMFI->addErrorMsg(
-                amd::CompilerErrorMessage[NO_IMAGE_SUPPORT]);
-            ++curReg;
-          }
-        } else {
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1);
-        }
-      } else if (CT->isStructTy()
-          && PT->getAddressSpace() == AMDILAS::PRIVATE_ADDRESS) {
-        StructType *ST = dyn_cast<StructType>(CT);
-        bool i1d  = ST->getName() == "struct._image1d_t";
-        bool i1da = ST->getName() == "struct._image1d_array_t";
-        bool i1db = ST->getName() == "struct._image1d_buffer_t";
-        bool i2d  = ST->getName() == "struct._image2d_t";
-        bool i2da = ST->getName() == "struct._image2d_array_t";
-        bool i3d  = ST->getName() == "struct._image3d_t";
-        bool is_image = i1d || i1da || i1db || i2d || i2da || i3d;
-        if (is_image) {
-          if (mSTM->device()->isSupported(AMDILDeviceInfo::Images)) {
-            printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer,
-                1, NULL, 0, is_image);
-          } else {
-            mMFI->addErrorMsg(amd::CompilerErrorMessage[NO_IMAGE_SUPPORT]);
-            ++curReg;
-          }
-        } else {
-          if (count) {
-            // Anytime we declare a literal, we need to reserve it, if it
-            // is not emitted in emitLiterals.
-            mMFI->addReservedLiterals(1);
-            O << "dcl_literal l" << mMFI->getNumLiterals() << ", "
-              << -stackSize << ", " << stackSize << ", 16, "
-              << stackOffset << "\n";
-          }
-          ++count;
-          size_t structSize;
-          structSize = (getTypeSize(ST) + 15) & ~15;
-          stackOffset += structSize;
-#ifdef UPSTREAM_LLVM
-          O << "mov " << RegNames->getRegisterName((curReg)) << ", l"
-            << mMFI->getNumLiterals()<< ".w\n";
-          if (!displaced1) {
-            O << "mov r1011, r1\n";
-            displaced1 = true;
-          }
-          if (!displaced2 && strcmp(RegNames->getRegisterName(curReg), "r1")) {
-            O << "mov r1010, r2\n";
-            displaced2 = true;
-          }
-#endif
-          printCopyStructPrivate(ST, O, structSize, Buffer, mMFI->getNumLiterals(),
-              Counter);
-          ++curReg;
-        }
-      } else if (CT->isIntOrIntVectorTy()
-          || CT->isFPOrFPVectorTy()
-          || CT->isArrayTy()
-          || CT->isPointerTy()
-          || PT->getAddressSpace() != AMDILAS::PRIVATE_ADDRESS) {
-        if (PT->getAddressSpace() == AMDILAS::CONSTANT_ADDRESS) {
-          const kernel& krnl = mGM->getKernel(F->getName());
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer,
-              1, NULL, 0, false, 
-              mGM->usesHWConstant(krnl, I->getName()));
-        } else if (PT->getAddressSpace() == AMDILAS::REGION_ADDRESS) {
-          // TODO: If we are region address space, the first region pointer, no
-          // array pointers exist, and hardware RegionMem is enabled then we can
-          // zero out register as the initial offset is zero.
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1);
-        } else if (PT->getAddressSpace() == AMDILAS::LOCAL_ADDRESS) {
-          // TODO: If we are local address space, the first local pointer, no
-          // array pointers exist, and hardware LocalMem is enabled then we can
-          // zero out register as the initial offset is zero.
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1);
-        } else {
-          printConstantToRegMapping(RegNames, curReg, O, Counter, Buffer, 1);
-        }
-      } else {
-        assert(0 && "Current type is not supported!");
-        mMFI->addErrorMsg(amd::CompilerErrorMessage[INTERNAL_ERROR]);
-        ++curReg;
-      }
-    } else {
-      assert(0 && "Current type is not supported!");
-      mMFI->addErrorMsg(amd::CompilerErrorMessage[INTERNAL_ERROR]);
-      ++curReg;
-    }
-  }
-  if (displaced1) {
-    O << "mov r1, r1011\n";
-  }
-  if (displaced2) {
-    O << "mov r2, r1010\n";
-  }
-  if (mSTM->device()->usesHardware(AMDILDeviceInfo::ConstantMem)) {
-    const kernel& krnl = mGM->getKernel(F->getName());
-    uint32_t constNum = 0;
-    for (uint32_t x = 0; x < mSTM->device()->getMaxNumCBs(); ++x) {
-      if (krnl.constSizes[x]) {
-        O << "dcl_cb cb" << x + CB_BASE_OFFSET;
-        O << "[" << (((krnl.constSizes[x] + 15) & ~15) >> 4) << "]\n";
-        ++constNum;
-        mMFI->setUsesMem(AMDILDevice::CONSTANT_ID);
-      }
-    }
-    // TODO: If we run out of constant resources, we need to push some of the
-    // constant pointers to the software emulated section.
-    if (constNum > mSTM->device()->getMaxNumCBs()) {
-      assert(0 && "Max constant buffer limit passed!");
-      mMFI->addErrorMsg(amd::CompilerErrorMessage[INSUFFICIENT_CONSTANT_RESOURCES]);
-    }
-  }
-}
-
-  const char *
-AMDILKernelManager::getTypeName(const Type *ptr, const char *symTab)
-{
-  // symTab argument is ignored...
-  LLVMContext& ctx = ptr->getContext();
-  switch (ptr->getTypeID()) {
-    case Type::StructTyID:
-      {
-        const StructType *ST = cast<StructType>(ptr);
-        if (!ST->isOpaque())
-          return "struct";
-        // ptr is a pre-LLVM 3.0 "opaque" type.
-        StringRef name = ST->getName();
-        if (name.equals( "struct._event_t" ))         return "event";
-        if (name.equals( "struct._image1d_t" ))       return "image1d";
-        if (name.equals( "struct._image1d_array_t" )) return "image1d_array";
-        if (name.equals( "struct._image2d_t" ))       return "image2d";
-        if (name.equals( "struct._image2d_array_t" )) return "image2d_array";
-        if (name.equals( "struct._image3d_t" ))       return "image3d";
-        if (name.equals( "struct._counter32_t" ))     return "counter32";
-        if (name.equals( "struct._counter64_t" ))     return "counter64";
-        return "opaque";
-        break;
-      }
-    case Type::FloatTyID:
-      return "float";
-    case Type::DoubleTyID: 
-      {
-        const AMDILSubtarget *mSTM= mTM->getSubtargetImpl();
-        if (!mSTM->device()->usesHardware(AMDILDeviceInfo::DoubleOps)) {
-          mMFI->addErrorMsg(amd::CompilerErrorMessage[DOUBLE_NOT_SUPPORTED]);
-        }
-        return "double";
-      }
-    case Type::IntegerTyID: 
-      {
-        if (ptr == Type::getInt8Ty(ctx)) {
-          return "i8";
-        } else if (ptr == Type::getInt16Ty(ctx)) {
-          return "i16";
-        } else if (ptr == Type::getInt32Ty(ctx)) {
-          return "i32";
-        } else if(ptr == Type::getInt64Ty(ctx)) {
-          return "i64";
-        }
-        break;
-      }
-    default:
-      break;
-    case Type::ArrayTyID: 
-      {
-        const ArrayType *AT = cast<ArrayType>(ptr);
-        const Type *name = AT->getElementType();
-        return getTypeName(name, symTab);
-        break;
-      }
-    case Type::VectorTyID: 
-      {
-        const VectorType *VT = cast<VectorType>(ptr);
-        const Type *name = VT->getElementType();
-        return getTypeName(name, symTab);
-        break;
-      }
-    case Type::PointerTyID: 
-      {
-        const PointerType *PT = cast<PointerType>(ptr);
-        const Type *name = PT->getElementType();
-        return getTypeName(name, symTab);
-        break;
-      }
-    case Type::FunctionTyID: 
-      {
-        const FunctionType *FT = cast<FunctionType>(ptr);
-        const Type *name = FT->getReturnType();
-        return getTypeName(name, symTab);
-        break;
-      }
-  }
-  ptr->dump();
-  mMFI->addErrorMsg(amd::CompilerErrorMessage[UNKNOWN_TYPE_NAME]);
-  return "unknown";
-}
-
-void AMDILKernelManager::emitLiterals(llvm::raw_ostream &O) {
-  char buffer[256];
-  std::map<uint32_t, uint32_t>::iterator ilb, ile;
-  for (ilb = mMFI->begin_32(), ile = mMFI->end_32(); ilb != ile; ++ilb) {
-    uint32_t a = ilb->first;
-    O << "dcl_literal l" <<ilb->second<< ", ";
-    sprintf(buffer, "0x%08x, 0x%08x, 0x%08x, 0x%08x", a, a, a, a);
-    O << buffer << "; f32:i32 " << ilb->first << "\n";
-  }
-  std::map<uint64_t, uint32_t>::iterator llb, lle;
-  for (llb = mMFI->begin_64(), lle = mMFI->end_64(); llb != lle; ++llb) {
-    uint32_t v[2];
-    uint64_t a = llb->first;
-    memcpy(v, &a, sizeof(uint64_t));
-    O << "dcl_literal l" <<llb->second<< ", ";
-    sprintf(buffer, "0x%08x, 0x%08x, 0x%08x, 0x%08x; f64:i64 ",
-        v[0], v[1], v[0], v[1]);
-    O << buffer << llb->first << "\n";
-  }
-  std::map<std::pair<uint64_t, uint64_t>, uint32_t>::iterator vlb, vle;
-  for (vlb = mMFI->begin_128(), vle = mMFI->end_128(); vlb != vle; ++vlb) {
-    uint32_t v[2][2];
-    uint64_t a = vlb->first.first;
-    uint64_t b = vlb->first.second;
-    memcpy(v[0], &a, sizeof(uint64_t));
-    memcpy(v[1], &b, sizeof(uint64_t));
-    O << "dcl_literal l" << vlb->second << ", ";
-    sprintf(buffer, "0x%08x, 0x%08x, 0x%08x, 0x%08x; f128:i128 ",
-        v[0][0], v[0][1], v[1][0], v[1][1]);
-    O << buffer << vlb->first.first << vlb->first.second << "\n";
-  }
-}
-
-// If the value is not known, then the uav is set, otherwise the mValueIDMap
-// is used.
-void AMDILKernelManager::setUAVID(const Value *value, uint32_t ID) {
-  if (value) {
-    mValueIDMap[value] = ID;
-  }
-}
-
-uint32_t AMDILKernelManager::getUAVID(const Value *value) {
-  if (mValueIDMap.find(value) != mValueIDMap.end()) {
-    return mValueIDMap[value];
-  }
-
-  if (mSTM->device()->getGeneration() <= AMDILDeviceInfo::HD6XXX) {
-    return mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID);
-  } else {
-    return mSTM->device()->getResourceID(AMDILDevice::RAW_UAV_ID);
-  }
-}
-
diff --git a/src/gallium/drivers/radeon/AMDILKernelManager.h b/src/gallium/drivers/radeon/AMDILKernelManager.h
deleted file mode 100644
index d5eb296cbf2..00000000000
--- a/src/gallium/drivers/radeon/AMDILKernelManager.h
+++ /dev/null
@@ -1,177 +0,0 @@
-//===-- AMDILKernelManager.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 the metadata/abi management for the
-// ASM printer. Handles the parsing and generation of the metadata
-// for each kernel and keeps track of its arguments.
-//
-//==-----------------------------------------------------------------------===//
-#ifndef _AMDILKERNELMANAGER_H_
-#define _AMDILKERNELMANAGER_H_
-#include "AMDIL.h"
-#include "AMDILDevice.h"
-#include "llvm/ADT/DenseMap.h"
-#include "llvm/ADT/DenseSet.h"
-#include "llvm/ADT/SmallVector.h"
-#include "llvm/ADT/StringMap.h"
-#include "llvm/ADT/ValueMap.h"
-#include "llvm/CodeGen/MachineBasicBlock.h"
-#include "llvm/Function.h"
-
-#include <map>
-#include <set>
-#include <string>
-
-#define IMAGETYPE_2D 0
-#define IMAGETYPE_3D 1
-#define RESERVED_LIT_COUNT 6
-
-namespace llvm {
-class AMDILGlobalManager;
-class AMDILSubtarget;
-class AMDILMachineFunctionInfo;
-class AMDILTargetMachine;
-class AMDILAsmPrinter;
-class StructType;
-class Value;
-class TypeSymbolTable;
-class MachineFunction;
-class MachineInstr;
-class ConstantFP;
-class PrintfInfo;
-
-
-class AMDILKernelManager {
-public:
-  typedef enum {
-    RELEASE_ONLY,
-    DEBUG_ONLY,
-    ALWAYS
-  } ErrorMsgEnum;
-  AMDILKernelManager(AMDILTargetMachine *TM, AMDILGlobalManager *GM);
-  virtual ~AMDILKernelManager();
-  
-  /// Clear the state of the KernelManager putting it in its most initial state.
-  void clear();
-  void setMF(MachineFunction *MF);
-
-  /// Process the specific kernel parsing out the parameter information for the
-  /// kernel.
-  void processArgMetadata(llvm::raw_ostream &O,
-                          uint32_t buf, bool kernel);
-
-
-  /// Prints the header for the kernel which includes the groupsize declaration
-  /// and calculation of the local/group/global id's.
-  void printHeader(AMDILAsmPrinter *AsmPrinter, llvm::raw_ostream &O,
-                   const std::string &name);
-
-  virtual void printDecls(AMDILAsmPrinter *AsmPrinter, llvm::raw_ostream &O);
-  virtual void printGroupSize(llvm::raw_ostream &O);
-
-  /// Copies the data from the runtime setup constant buffers into registers so
-  /// that the program can correctly access memory or data that was set by the
-  /// host program.
-  void printArgCopies(llvm::raw_ostream &O, AMDILAsmPrinter* RegNames);
-
-  /// Prints out the end of the function.
-  void printFooter(llvm::raw_ostream &O);
-  
-  /// Prints out the metadata for the specific function depending if it is a
-  /// kernel or not.
-  void printMetaData(llvm::raw_ostream &O, uint32_t id, bool isKernel = false);
-  
-  /// Set bool value on whether to consider the function a kernel or a normal
-  /// function.
-  void setKernel(bool kernel);
-
-  /// Set the unique ID of the kernel/function.
-  void setID(uint32_t id);
-
-  /// Set the name of the kernel/function.
-  void setName(const std::string &name);
-
-  /// Flag to specify whether the function is a kernel or not.
-  bool isKernel();
-
-  /// Flag that specifies whether this function has a kernel wrapper.
-  bool wasKernel();
-
-  void getIntrinsicSetup(AMDILAsmPrinter *AsmPrinter, llvm::raw_ostream &O); 
-
-  // Returns whether a compiler needs to insert a write to memory or not.
-  bool useCompilerWrite(const MachineInstr *MI);
-
-  // Set the flag that there exists an image write.
-  void setImageWrite();
-  void setOutputInst();
-
-  const char *getTypeName(const Type *name, const char * symTab);
-
-  void emitLiterals(llvm::raw_ostream &O);
-
-  // Set the uav id for the specific pointer value.  If value is NULL, then the
-  // ID sets the default ID.
-  void setUAVID(const Value *value, uint32_t ID);
-
-  // Get the UAV id for the specific pointer value.
-  uint32_t getUAVID(const Value *value);
-
-private:
-
-  /// Helper function that prints the actual metadata and should only be called
-  /// by printMetaData.
-  void printKernelArgs(llvm::raw_ostream &O);
-  void printCopyStructPrivate(const StructType *ST,
-                              llvm::raw_ostream &O,
-                              size_t stackSize,
-                              uint32_t Buffer,
-                              uint32_t mLitIdx,
-                              uint32_t &counter);
-  virtual void
-  printConstantToRegMapping(AMDILAsmPrinter *RegNames,
-                            uint32_t &LII,
-                            llvm::raw_ostream &O,
-                            uint32_t &counter,
-                            uint32_t Buffer,
-                            uint32_t n,
-                            const char *lit = NULL,
-                            uint32_t fcall = 0,
-                            bool isImage = false,
-                            bool isHWCB = false);
-  void updatePtrArg(llvm::Function::const_arg_iterator Ip,
-                    int numWriteImages,
-                    int raw_uav_buffer,
-                    int counter,
-                    bool isKernel,
-                    const Function *F);
-  /// Name of the current kernel.
-  std::string mName;
-  uint32_t mUniqueID;
-  bool mIsKernel;
-  bool mWasKernel;
-  bool mCompilerWrite;
-  /// Flag to specify if an image write has occured or not in order to not add a
-  /// compiler specific write if no other writes to memory occured.
-  bool mHasImageWrite;
-  bool mHasOutputInst;
-  
-  /// Map from const Value * to UAV ID.
-  std::map<const Value *, uint32_t> mValueIDMap;
-
-  AMDILTargetMachine * mTM;
-  const AMDILSubtarget * mSTM;
-  AMDILGlobalManager * mGM;
-  /// This is the global offset of the printf string id's.
-  MachineFunction *mMF;
-  AMDILMachineFunctionInfo *mMFI;
-}; // class AMDILKernelManager
-
-} // llvm namespace
-#endif // _AMDILKERNELMANAGER_H_
diff --git a/src/gallium/drivers/radeon/AMDILLiteralManager.cpp b/src/gallium/drivers/radeon/AMDILLiteralManager.cpp
index 43167f57001..52bf7674e7d 100644
--- a/src/gallium/drivers/radeon/AMDILLiteralManager.cpp
+++ b/src/gallium/drivers/radeon/AMDILLiteralManager.cpp
@@ -12,7 +12,6 @@
 #include "AMDIL.h"
 
 #include "AMDILAlgorithms.tpp"
-#include "AMDILKernelManager.h"
 #include "AMDILMachineFunctionInfo.h"
 #include "AMDILSubtarget.h"
 #include "AMDILTargetMachine.h"
@@ -43,7 +42,6 @@ namespace {
     bool trackLiterals(MachineBasicBlock::iterator *bbb);
     TargetMachine &TM;
     const AMDILSubtarget *mSTM;
-    AMDILKernelManager *mKM;
     AMDILMachineFunctionInfo *mMFI;
     int32_t mLitIdx;
     bool mChanged;
@@ -71,7 +69,6 @@ bool AMDILLiteralManager::runOnMachineFunction(MachineFunction &MF) {
   const AMDILTargetMachine *amdtm =
     reinterpret_cast<const AMDILTargetMachine *>(&TM);
   mSTM = dynamic_cast<const AMDILSubtarget *>(amdtm->getSubtargetImpl());
-  mKM = const_cast<AMDILKernelManager *>(mSTM->getKernelManager());
   safeNestedForEach(MF.begin(), MF.end(), MF.begin()->begin(),
       std::bind1st(std::mem_fun(&AMDILLiteralManager::trackLiterals), this));
   return mChanged;
diff --git a/src/gallium/drivers/radeon/AMDILPeepholeOptimizer.cpp b/src/gallium/drivers/radeon/AMDILPeepholeOptimizer.cpp
index 9383bfcb77b..4859fe9df51 100644
--- a/src/gallium/drivers/radeon/AMDILPeepholeOptimizer.cpp
+++ b/src/gallium/drivers/radeon/AMDILPeepholeOptimizer.cpp
@@ -16,8 +16,6 @@
 
 #include "AMDILAlgorithms.tpp"
 #include "AMDILDevices.h"
-#include "AMDILGlobalManager.h"
-#include "AMDILKernelManager.h"
 #include "AMDILMachineFunctionInfo.h"
 #include "AMDILUtilityFunctions.h"
 #include "llvm/ADT/Statistic.h"
@@ -39,7 +37,6 @@ STATISTIC(PointerAssignments, "Number of dynamic pointer "
     "assigments discovered");
 STATISTIC(PointerSubtract, "Number of pointer subtractions discovered");
 #endif
-STATISTIC(LocalFuncs, "Number of get_local_size(N) functions removed");
 
 using namespace llvm;
 // The Peephole optimization pass is used to do simple last minute optimizations
@@ -76,7 +73,6 @@ private:
   void doIsConstCallConversionIfNeeded();
   bool mChanged;
   bool mDebug;
-  bool mRWGOpt;
   bool mConvertAtomics;
   CodeGenOpt::Level optLevel;
   // Run a series of tests to see if we can optimize a CALL instruction.
@@ -104,7 +100,6 @@ private:
   // specified then the result of get_local_size is known at compile time and
   // can be returned accordingly.
   bool isRWGLocalOpt(CallInst *CI);
-  void expandRWGLocalOpt(CallInst *CI);
   // On northern island cards, the division is slightly less accurate than on
   // previous generations, so we need to utilize a more accurate division. So we
   // can translate the accurate divide to a normal divide on all other cards.
@@ -251,15 +246,12 @@ AMDILPeepholeOpt::doAtomicConversionIfNeeded(Function &F)
   // arena path.
   Function::arg_iterator argB = F.arg_begin();
   Function::arg_iterator argE = F.arg_end();
-  AMDILKernelManager *KM = mSTM->getKernelManager();
   AMDILMachineFunctionInfo *mMFI = getAnalysis<MachineFunctionAnalysis>().getMF()
     .getInfo<AMDILMachineFunctionInfo>();
   for (; argB != argE; ++argB) {
     if (mSTM->device()->isSupported(AMDILDeviceInfo::ArenaUAV)) {
-      KM->setUAVID(argB,mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID));
       mMFI->uav_insert(mSTM->device()->getResourceID(AMDILDevice::ARENA_UAV_ID));
     } else {
-      KM->setUAVID(argB,mSTM->device()->getResourceID(AMDILDevice::GLOBAL_ID));
       mMFI->uav_insert(mSTM->device()->getResourceID(AMDILDevice::GLOBAL_ID));
     }
   }
@@ -276,16 +268,6 @@ AMDILPeepholeOpt::runOnFunction(Function &MF)
   }
   mCTX = &MF.getType()->getContext();
   mConvertAtomics = true;
-  if (dumpAllIntoArena(MF)) {
-    for (Function::const_arg_iterator cab = MF.arg_begin(),
-         cae = MF.arg_end(); cab != cae; ++cab) {
-      const Argument *arg = cab;
-      AMDILKernelManager *KM = mSTM->getKernelManager();
-      KM->setUAVID(getBasePointerValue(arg),
-          mSTM->device()->getResourceID(AMDILDevice::GLOBAL_ID));
-    }
-  }
-  mRWGOpt = mSTM->getGlobalManager()->hasRWG(MF.getName());
   safeNestedForEach(MF.begin(), MF.end(), MF.begin()->begin(),
      std::bind1st(std::mem_fun(&AMDILPeepholeOpt::instLevelOptimizations),
                   this));
@@ -313,10 +295,6 @@ AMDILPeepholeOpt::optimizeCallInst(BasicBlock::iterator *bbb)
     CI->eraseFromParent();
     return true;
   }
-  if (isRWGLocalOpt(CI)) {
-    expandRWGLocalOpt(CI);
-    return false;
-  }
   if (propagateSamplerInst(CI)) {
     return false;
   }
@@ -390,26 +368,7 @@ AMDILPeepholeOpt::optimizeCallInst(BasicBlock::iterator *bbb)
   }
   StringRef name = F->getName();
   if (name.startswith("__atom") && name.find("_g") != StringRef::npos) {
-    Value *ptr = CI->getOperand(0);
-    const Value *basePtr = getBasePointerValue(ptr);
-    const Argument *Arg = dyn_cast<Argument>(basePtr);
-    if (Arg) {
-      AMDILGlobalManager *GM = mSTM->getGlobalManager();
-      int32_t id = GM->getArgID(Arg);
-      if (id >= 0) {
-        std::stringstream ss;
-        ss << name.data() << "_" << id << '\n';
-        std::string val;
-        ss >> val;
-        F = dyn_cast<Function>(
-              F->getParent() ->getOrInsertFunction(val, F->getFunctionType()));
-        atomicFuncs.push_back(std::make_pair <CallInst*, Function*>(CI, F));
-      } else {
-        mConvertAtomics = false;
-      }
-    } else {
-      mConvertAtomics = false;
-    }
+    mConvertAtomics = false;
   }
   return false;
 }
@@ -1088,28 +1047,11 @@ AMDILPeepholeOpt::expandSigned24BitOps(CallInst *CI)
 bool 
 AMDILPeepholeOpt::isRWGLocalOpt(CallInst *CI) 
 {
-  return (CI != NULL && mRWGOpt
+  return (CI != NULL
           && CI->getOperand(CI->getNumOperands() - 1)->getName() 
           == "__amdil_get_local_size_int");
 }
 
-void 
-AMDILPeepholeOpt::expandRWGLocalOpt(CallInst *CI) 
-{
-  assert(isRWGLocalOpt(CI) &&
-         "This optmization only works when the call inst is get_local_size!");
-  std::vector<Constant *> consts;
-  for (uint32_t x = 0; x < 3; ++x) {
-    uint32_t val = mSTM->getGlobalManager()->getLocal(mF->getName(), x);
-    consts.push_back(ConstantInt::get(Type::getInt32Ty(*mCTX), val));
-  }
-  consts.push_back(ConstantInt::get(Type::getInt32Ty(*mCTX), 0));
-  Value *cVec = ConstantVector::get(consts);
-  CI->replaceAllUsesWith(cVec);
-  ++LocalFuncs;
-  return;
-}
-
 bool 
 AMDILPeepholeOpt::convertAccurateDivide(CallInst *CI) 
 {
diff --git a/src/gallium/drivers/radeon/AMDILPrintfConvert.cpp b/src/gallium/drivers/radeon/AMDILPrintfConvert.cpp
index 95614f477c0..17d9e8e3013 100644
--- a/src/gallium/drivers/radeon/AMDILPrintfConvert.cpp
+++ b/src/gallium/drivers/radeon/AMDILPrintfConvert.cpp
@@ -15,7 +15,6 @@
 #endif
 
 #include "AMDILAlgorithms.tpp"
-#include "AMDILKernelManager.h"
 #include "AMDILMachineFunctionInfo.h"
 #include "AMDILModuleInfo.h"
 #include "AMDILTargetMachine.h"
@@ -50,7 +49,6 @@ namespace
         private:
             bool expandPrintf(BasicBlock::iterator *bbb);
             AMDILMachineFunctionInfo *mMFI;
-            AMDILKernelManager *mKM;
             bool mChanged;
             SmallVector<int64_t, DEFAULT_VEC_SLOTS> bVecMap;
     };
@@ -256,7 +254,6 @@ AMDILPrintfConvert::expandPrintf(BasicBlock::iterator *bbb)
 AMDILPrintfConvert::runOnFunction(Function &MF)
 {
     mChanged = false;
-    mKM = TM.getSubtarget<AMDILSubtarget>().getKernelManager();
     mMFI = getAnalysis<MachineFunctionAnalysis>().getMF()
           .getInfo<AMDILMachineFunctionInfo>();
     bVecMap.clear();
diff --git a/src/gallium/drivers/radeon/AMDILSubtarget.cpp b/src/gallium/drivers/radeon/AMDILSubtarget.cpp
index 898833d9c0e..11b6bbe0c01 100644
--- a/src/gallium/drivers/radeon/AMDILSubtarget.cpp
+++ b/src/gallium/drivers/radeon/AMDILSubtarget.cpp
@@ -14,8 +14,6 @@
 #include "AMDILSubtarget.h"
 #include "AMDIL.h"
 #include "AMDILDevices.h"
-#include "AMDILGlobalManager.h"
-#include "AMDILKernelManager.h"
 #include "AMDILUtilityFunctions.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringExtras.h"
diff --git a/src/gallium/drivers/radeon/Makefile.sources b/src/gallium/drivers/radeon/Makefile.sources
index bfffdd1eeb8..524289f5e0d 100644
--- a/src/gallium/drivers/radeon/Makefile.sources
+++ b/src/gallium/drivers/radeon/Makefile.sources
@@ -26,13 +26,11 @@ CPP_SOURCES := \
 	AMDILEvergreenDevice.cpp	\
 	AMDILELFWriterInfo.cpp		\
 	AMDILFrameLowering.cpp		\
-	AMDILGlobalManager.cpp		\
 	AMDILInliner.cpp		\
 	AMDILInstrInfo.cpp		\
 	AMDILIntrinsicInfo.cpp		\
 	AMDILISelDAGToDAG.cpp		\
 	AMDILISelLowering.cpp		\
-	AMDILKernelManager.cpp		\
 	AMDILLiteralManager.cpp		\
 	AMDILMachineFunctionInfo.cpp	\
 	AMDILMachinePeephole.cpp	\