diff options
Diffstat (limited to 'lib/Target/NVPTX/NVPTXAsmPrinter.cpp')
-rw-r--r-- | lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 211 |
1 files changed, 107 insertions, 104 deletions
diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index ce5d78a..ff73931 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -68,11 +68,12 @@ InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, namespace { /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V /// depends. -void DiscoverDependentGlobals(Value *V, DenseSet<GlobalVariable *> &Globals) { - if (GlobalVariable *GV = dyn_cast<GlobalVariable>(V)) +void DiscoverDependentGlobals(const Value *V, + DenseSet<const GlobalVariable *> &Globals) { + if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V)) Globals.insert(GV); else { - if (User *U = dyn_cast<User>(V)) { + if (const User *U = dyn_cast<User>(V)) { for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) { DiscoverDependentGlobals(U->getOperand(i), Globals); } @@ -84,8 +85,9 @@ void DiscoverDependentGlobals(Value *V, DenseSet<GlobalVariable *> &Globals) { /// instances to be emitted, but only after any dependents have been added /// first. void VisitGlobalVariableForEmission( - GlobalVariable *GV, SmallVectorImpl<GlobalVariable *> &Order, - DenseSet<GlobalVariable *> &Visited, DenseSet<GlobalVariable *> &Visiting) { + const GlobalVariable *GV, SmallVectorImpl<const GlobalVariable *> &Order, + DenseSet<const GlobalVariable *> &Visited, + DenseSet<const GlobalVariable *> &Visiting) { // Have we already visited this one? if (Visited.count(GV)) return; @@ -98,12 +100,12 @@ void VisitGlobalVariableForEmission( Visiting.insert(GV); // Make sure we visit all dependents first - DenseSet<GlobalVariable *> Others; + DenseSet<const GlobalVariable *> Others; for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i) DiscoverDependentGlobals(GV->getOperand(i), Others); - for (DenseSet<GlobalVariable *>::iterator I = Others.begin(), - E = Others.end(); + for (DenseSet<const GlobalVariable *>::iterator I = Others.begin(), + E = Others.end(); I != E; ++I) VisitGlobalVariableForEmission(*I, Order, Visited, Visiting); @@ -405,6 +407,11 @@ void NVPTXAsmPrinter::EmitFunctionEntryLabel() { SmallString<128> Str; raw_svector_ostream O(Str); + if (!GlobalsEmitted) { + emitGlobals(*MF->getFunction()->getParent()); + GlobalsEmitted = true; + } + // Set up MRI = &MF->getRegInfo(); F = MF->getFunction(); @@ -429,9 +436,7 @@ void NVPTXAsmPrinter::EmitFunctionEntryLabel() { } void NVPTXAsmPrinter::EmitFunctionBodyStart() { - const TargetRegisterInfo &TRI = *TM.getRegisterInfo(); - unsigned numRegClasses = TRI.getNumRegClasses(); - VRidGlobal2LocalMap = new std::map<unsigned, unsigned>[numRegClasses + 1]; + VRegMapping.clear(); OutStreamer.EmitRawText(StringRef("{\n")); setAndEmitFunctionVirtualRegisters(*MF); @@ -443,7 +448,7 @@ void NVPTXAsmPrinter::EmitFunctionBodyStart() { void NVPTXAsmPrinter::EmitFunctionBodyEnd() { OutStreamer.EmitRawText(StringRef("}\n")); - delete[] VRidGlobal2LocalMap; + VRegMapping.clear(); } void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, @@ -500,9 +505,8 @@ void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F, void NVPTXAsmPrinter::getVirtualRegisterName(unsigned vr, bool isVec, raw_ostream &O) { const TargetRegisterClass *RC = MRI->getRegClass(vr); - unsigned id = RC->getID(); - std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[id]; + DenseMap<unsigned, unsigned> ®map = VRegMapping[RC]; unsigned mapped_vr = regmap[vr]; if (!isVec) { @@ -695,7 +699,7 @@ void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) { else O << ".func "; printReturnValStr(F, O); - O << *CurrentFnSym << "\n"; + O << *Mang->getSymbol(F) << "\n"; emitFunctionParamList(F, O); O << ";\n"; } @@ -795,7 +799,7 @@ static bool useFuncSeen(const Constant *C, return false; } -void NVPTXAsmPrinter::emitDeclarations(Module &M, raw_ostream &O) { +void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) { llvm::DenseMap<const Function *, bool> seenMap; for (Module::const_iterator FI = M.begin(), FE = M.end(); FI != FE; ++FI) { const Function *F = FI; @@ -805,7 +809,6 @@ void NVPTXAsmPrinter::emitDeclarations(Module &M, raw_ostream &O) { continue; if (F->getIntrinsicID()) continue; - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); continue; } @@ -817,14 +820,12 @@ void NVPTXAsmPrinter::emitDeclarations(Module &M, raw_ostream &O) { // The use is in the initialization of a global variable // that is a function pointer, so print a declaration // for the original function - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); break; } // Emit a declaration of this function if the function that // uses this constant expr has already been seen. if (useFuncSeen(C, seenMap)) { - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); break; } @@ -844,7 +845,6 @@ void NVPTXAsmPrinter::emitDeclarations(Module &M, raw_ostream &O) { // appearing in the module before the callee. so print out // a declaration for the callee. if (seenMap.find(caller) != seenMap.end()) { - CurrentFnSym = Mang->getSymbol(F); emitDeclaration(F, O); break; } @@ -909,7 +909,7 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { const_cast<TargetLoweringObjectFile &>(getObjFileLowering()) .Initialize(OutContext, TM); - Mang = new Mangler(OutContext, *TM.getDataLayout()); + Mang = new Mangler(OutContext, &TM); // Emit header before any dwarf directives are emitted below. emitHeader(M, OS1); @@ -921,6 +921,12 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) recordAndEmitFilenames(M); + GlobalsEmitted = false; + + return false; // success +} + +void NVPTXAsmPrinter::emitGlobals(const Module &M) { SmallString<128> Str2; raw_svector_ostream OS2(Str2); @@ -931,13 +937,13 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { // global variable in order, and ensure that we emit it *after* its dependent // globals. We use a little extra memory maintaining both a set and a list to // have fast searches while maintaining a strict ordering. - SmallVector<GlobalVariable *, 8> Globals; - DenseSet<GlobalVariable *> GVVisited; - DenseSet<GlobalVariable *> GVVisiting; + SmallVector<const GlobalVariable *, 8> Globals; + DenseSet<const GlobalVariable *> GVVisited; + DenseSet<const GlobalVariable *> GVVisiting; // Visit each global variable, in order - for (Module::global_iterator I = M.global_begin(), E = M.global_end(); I != E; - ++I) + for (Module::const_global_iterator I = M.global_begin(), E = M.global_end(); + I != E; ++I) VisitGlobalVariableForEmission(I, Globals, GVVisited, GVVisiting); assert(GVVisited.size() == M.getGlobalList().size() && @@ -951,7 +957,6 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) { OS2 << '\n'; OutStreamer.EmitRawText(OS2.str()); - return false; // success } void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O) { @@ -989,6 +994,14 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O) { } bool NVPTXAsmPrinter::doFinalization(Module &M) { + + // If we did not emit any functions, then the global declarations have not + // yet been emitted. + if (!GlobalsEmitted) { + emitGlobals(M); + GlobalsEmitted = true; + } + // XXX Temproarily remove global variables so that doFinalization() will not // emit them again (global variables are emitted at beginning). @@ -1063,7 +1076,8 @@ void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V, } } -void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable *GVar, raw_ostream &O, +void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, + raw_ostream &O, bool processDemoted) { // Skip meta data @@ -1107,10 +1121,10 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable *GVar, raw_ostream &O, if (llvm::isSampler(*GVar)) { O << ".global .samplerref " << llvm::getSamplerName(*GVar); - Constant *Initializer = NULL; + const Constant *Initializer = NULL; if (GVar->hasInitializer()) Initializer = GVar->getInitializer(); - ConstantInt *CI = NULL; + const ConstantInt *CI = NULL; if (Initializer) CI = dyn_cast<ConstantInt>(Initializer); if (CI) { @@ -1183,7 +1197,7 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable *GVar, raw_ostream &O, if (localDecls.find(demotedFunc) != localDecls.end()) localDecls[demotedFunc].push_back(GVar); else { - std::vector<GlobalVariable *> temp; + std::vector<const GlobalVariable *> temp; temp.push_back(GVar); localDecls[demotedFunc] = temp; } @@ -1199,17 +1213,20 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable *GVar, raw_ostream &O, if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa<PointerType>(ETy)) { O << " ."; - O << getPTXFundamentalTypeStr(ETy, false); + // Special case: ABI requires that we use .u8 for predicates + if (ETy->isIntegerTy(1)) + O << "u8"; + else + O << getPTXFundamentalTypeStr(ETy, false); O << " "; O << *Mang->getSymbol(GVar); // Ptx allows variable initilization only for constant and global state // spaces. if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || - (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) || (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) && GVar->hasInitializer()) { - Constant *Initializer = GVar->getInitializer(); + const Constant *Initializer = GVar->getInitializer(); if (!Initializer->isNullValue()) { O << " = "; printScalarConstant(Initializer, O); @@ -1230,10 +1247,9 @@ void NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable *GVar, raw_ostream &O, // Ptx allows variable initilization only for constant and // global state spaces. if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || - (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) || (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) && GVar->hasInitializer()) { - Constant *Initializer = GVar->getInitializer(); + const Constant *Initializer = GVar->getInitializer(); if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) { AggBuffer aggBuffer(ElementSize, O, *this); bufferAggregateConstant(Initializer, &aggBuffer); @@ -1283,7 +1299,7 @@ void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) { if (localDecls.find(f) == localDecls.end()) return; - std::vector<GlobalVariable *> &gvars = localDecls[f]; + std::vector<const GlobalVariable *> &gvars = localDecls[f]; for (unsigned i = 0, e = gvars.size(); i != e; ++i) { O << "\t// demoted variable\n\t"; @@ -1301,14 +1317,6 @@ void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace, O << "global"; break; case llvm::ADDRESS_SPACE_CONST: - // This logic should be consistent with that in - // getCodeAddrSpace() (NVPTXISelDATToDAT.cpp) - if (nvptxSubtarget.hasGenericLdSt()) - O << "global"; - else - O << "const"; - break; - case llvm::ADDRESS_SPACE_CONST_NOT_GEN: O << "const"; break; case llvm::ADDRESS_SPACE_SHARED: @@ -1448,7 +1456,7 @@ void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I, int paramIndex, raw_ostream &O) { if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) - O << *CurrentFnSym << "_param_" << paramIndex; + O << *Mang->getSymbol(I->getParent()) << "_param_" << paramIndex; else { std::string argName = I->getName(); const char *p = argName.c_str(); @@ -1507,11 +1515,13 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { if (llvm::isImage(*I)) { std::string sname = I->getName(); if (llvm::isImageWriteOnly(*I)) - O << "\t.param .surfref " << *CurrentFnSym << "_param_" << paramIndex; + O << "\t.param .surfref " << *Mang->getSymbol(F) << "_param_" + << paramIndex; else // Default image is read_only - O << "\t.param .texref " << *CurrentFnSym << "_param_" << paramIndex; + O << "\t.param .texref " << *Mang->getSymbol(F) << "_param_" + << paramIndex; } else // Should be llvm::isSampler(*I) - O << "\t.param .samplerref " << *CurrentFnSym << "_param_" + O << "\t.param .samplerref " << *Mang->getSymbol(F) << "_param_" << paramIndex; continue; } @@ -1546,14 +1556,13 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { default: O << ".ptr "; break; - case llvm::ADDRESS_SPACE_CONST_NOT_GEN: + case llvm::ADDRESS_SPACE_CONST: O << ".ptr .const "; break; case llvm::ADDRESS_SPACE_SHARED: O << ".ptr .shared "; break; case llvm::ADDRESS_SPACE_GLOBAL: - case llvm::ADDRESS_SPACE_CONST: O << ".ptr .global "; break; } @@ -1564,7 +1573,13 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) { } // non-pointer scalar to kernel func - O << "\t.param ." << getPTXFundamentalTypeStr(Ty) << " "; + O << "\t.param ."; + // Special case: predicate operands become .u8 types + if (Ty->isIntegerTy(1)) + O << "u8"; + else + O << getPTXFundamentalTypeStr(Ty); + O << " "; printParamName(I, paramIndex, O); continue; } @@ -1680,48 +1695,36 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters( for (unsigned i = 0; i < numVRs; i++) { unsigned int vr = TRI->index2VirtReg(i); const TargetRegisterClass *RC = MRI->getRegClass(vr); - std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[RC->getID()]; + DenseMap<unsigned, unsigned> ®map = VRegMapping[RC]; int n = regmap.size(); regmap.insert(std::make_pair(vr, n + 1)); } // Emit register declarations // @TODO: Extract out the real register usage - O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .s64 %rl<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n"; - O << "\t.reg .f64 %fl<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .s64 %rl<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n"; + // O << "\t.reg .f64 %fl<" << NVPTXNumRegisters << ">;\n"; // Emit declaration of the virtual registers or 'physical' registers for // each register class - //for (unsigned i=0; i< numRegClasses; i++) { - // std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[i]; - // const TargetRegisterClass *RC = TRI->getRegClass(i); - // std::string rcname = getNVPTXRegClassName(RC); - // std::string rcStr = getNVPTXRegClassStr(RC); - // //int n = regmap.size(); - // if (!isNVPTXVectorRegClass(RC)) { - // O << "\t.reg " << rcname << " \t" << rcStr << "<" - // << NVPTXNumRegisters << ">;\n"; - // } - - // Only declare those registers that may be used. And do not emit vector - // registers as - // they are all elementized to scalar registers. - //if (n && !isNVPTXVectorRegClass(RC)) { - // if (RegAllocNilUsed) { - // O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1) - // << ">;\n"; - // } - // else { - // O << "\t.reg " << rcname << " \t" << StrToUpper(rcStr) - // << "<" << 32 << ">;\n"; - // } - //} - //} + for (unsigned i=0; i< TRI->getNumRegClasses(); i++) { + const TargetRegisterClass *RC = TRI->getRegClass(i); + DenseMap<unsigned, unsigned> ®map = VRegMapping[RC]; + std::string rcname = getNVPTXRegClassName(RC); + std::string rcStr = getNVPTXRegClassStr(RC); + int n = regmap.size(); + + // Only declare those registers that may be used. + if (n) { + O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1) + << ">;\n"; + } + } OutStreamer.EmitRawText(O.str()); } @@ -1751,12 +1754,12 @@ void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) { O << utohexstr(API.getZExtValue()); } -void NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) { - if (ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) { +void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) { + if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) { O << CI->getValue(); return; } - if (ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) { + if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) { printFPConstant(CFP, O); return; } @@ -1764,13 +1767,13 @@ void NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) { O << "0"; return; } - if (GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) { + if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) { O << *Mang->getSymbol(GVar); return; } - if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { - Value *v = Cexpr->stripPointerCasts(); - if (GlobalValue *GVar = dyn_cast<GlobalValue>(v)) { + if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { + const Value *v = Cexpr->stripPointerCasts(); + if (const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) { O << *Mang->getSymbol(GVar); return; } else { @@ -1781,7 +1784,7 @@ void NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) { llvm_unreachable("Not scalar type found in printScalarConstant()"); } -void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, +void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes, AggBuffer *aggBuffer) { const DataLayout *TD = TM.getDataLayout(); @@ -1809,13 +1812,13 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, ptr = (unsigned char *)&int16; aggBuffer->addBytes(ptr, 2, Bytes); } else if (ETy == Type::getInt32Ty(CPV->getContext())) { - if (ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) { + if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) { int int32 = (int)(constInt->getZExtValue()); ptr = (unsigned char *)&int32; aggBuffer->addBytes(ptr, 4, Bytes); break; - } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { - if (ConstantInt *constInt = dyn_cast<ConstantInt>( + } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { + if (const ConstantInt *constInt = dyn_cast<ConstantInt>( ConstantFoldConstantExpression(Cexpr, TD))) { int int32 = (int)(constInt->getZExtValue()); ptr = (unsigned char *)&int32; @@ -1831,13 +1834,13 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, } llvm_unreachable("unsupported integer const type"); } else if (ETy == Type::getInt64Ty(CPV->getContext())) { - if (ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) { + if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) { long long int64 = (long long)(constInt->getZExtValue()); ptr = (unsigned char *)&int64; aggBuffer->addBytes(ptr, 8, Bytes); break; - } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { - if (ConstantInt *constInt = dyn_cast<ConstantInt>( + } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { + if (const ConstantInt *constInt = dyn_cast<ConstantInt>( ConstantFoldConstantExpression(Cexpr, TD))) { long long int64 = (long long)(constInt->getZExtValue()); ptr = (unsigned char *)&int64; @@ -1858,7 +1861,7 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, } case Type::FloatTyID: case Type::DoubleTyID: { - ConstantFP *CFP = dyn_cast<ConstantFP>(CPV); + const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV); const Type *Ty = CFP->getType(); if (Ty == Type::getFloatTy(CPV->getContext())) { float float32 = (float) CFP->getValueAPF().convertToFloat(); @@ -1874,10 +1877,10 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, break; } case Type::PointerTyID: { - if (GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) { + if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) { aggBuffer->addSymbol(GVar); - } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { - Value *v = Cexpr->stripPointerCasts(); + } else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { + const Value *v = Cexpr->stripPointerCasts(); aggBuffer->addSymbol(v); } unsigned int s = TD->getTypeAllocSize(CPV->getType()); @@ -1906,7 +1909,7 @@ void NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, } } -void NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV, +void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV, AggBuffer *aggBuffer) { const DataLayout *TD = TM.getDataLayout(); int Bytes; |