1 //===-- PTXAsmPrinter.cpp - PTX LLVM assembly writer ----------------------===// 2 // 3 // The LLVM Compiler Infrastructure 4 // 5 // This file is distributed under the University of Illinois Open Source 6 // License. See LICENSE.TXT for details. 7 // 8 //===----------------------------------------------------------------------===// 9 // 10 // This file contains a printer that converts from our internal representation 11 // of machine-dependent LLVM code to PTX assembly language. 12 // 13 //===----------------------------------------------------------------------===// 14 15 #define DEBUG_TYPE "ptx-asm-printer" 16 17 #include "PTX.h" 18 #include "PTXAsmPrinter.h" 19 #include "PTXMachineFunctionInfo.h" 20 #include "PTXParamManager.h" 21 #include "PTXRegisterInfo.h" 22 #include "PTXTargetMachine.h" 23 #include "llvm/Argument.h" 24 #include "llvm/DerivedTypes.h" 25 #include "llvm/Function.h" 26 #include "llvm/Module.h" 27 #include "llvm/ADT/SmallString.h" 28 #include "llvm/ADT/StringExtras.h" 29 #include "llvm/ADT/Twine.h" 30 #include "llvm/Analysis/DebugInfo.h" 31 #include "llvm/CodeGen/AsmPrinter.h" 32 #include "llvm/CodeGen/MachineFrameInfo.h" 33 #include "llvm/CodeGen/MachineInstr.h" 34 #include "llvm/CodeGen/MachineRegisterInfo.h" 35 #include "llvm/MC/MCContext.h" 36 #include "llvm/MC/MCExpr.h" 37 #include "llvm/MC/MCInst.h" 38 #include "llvm/MC/MCStreamer.h" 39 #include "llvm/MC/MCSymbol.h" 40 #include "llvm/Target/Mangler.h" 41 #include "llvm/Target/TargetLoweringObjectFile.h" 42 #include "llvm/Support/CommandLine.h" 43 #include "llvm/Support/Debug.h" 44 #include "llvm/Support/ErrorHandling.h" 45 #include "llvm/Support/MathExtras.h" 46 #include "llvm/Support/Path.h" 47 #include "llvm/Support/TargetRegistry.h" 48 #include "llvm/Support/raw_ostream.h" 49 50 using namespace llvm; 51 52 static const char PARAM_PREFIX[] = "__param_"; 53 static const char RETURN_PREFIX[] = "__ret_"; 54 55 static const char *getRegisterTypeName(unsigned RegNo, 56 const MachineRegisterInfo& MRI) { 57 const TargetRegisterClass *TRC = MRI.getRegClass(RegNo); 58 59 #define TEST_REGCLS(cls, clsstr) \ 60 if (PTX::cls ## RegisterClass == TRC) return # clsstr; 61 62 TEST_REGCLS(RegPred, pred); 63 TEST_REGCLS(RegI16, b16); 64 TEST_REGCLS(RegI32, b32); 65 TEST_REGCLS(RegI64, b64); 66 TEST_REGCLS(RegF32, b32); 67 TEST_REGCLS(RegF64, b64); 68 #undef TEST_REGCLS 69 70 llvm_unreachable("Not in any register class!"); 71 return NULL; 72 } 73 74 static const char *getStateSpaceName(unsigned addressSpace) { 75 switch (addressSpace) { 76 default: llvm_unreachable("Unknown state space"); 77 case PTXStateSpace::Global: return "global"; 78 case PTXStateSpace::Constant: return "const"; 79 case PTXStateSpace::Local: return "local"; 80 case PTXStateSpace::Parameter: return "param"; 81 case PTXStateSpace::Shared: return "shared"; 82 } 83 return NULL; 84 } 85 86 static const char *getTypeName(Type* type) { 87 while (true) { 88 switch (type->getTypeID()) { 89 default: llvm_unreachable("Unknown type"); 90 case Type::FloatTyID: return ".f32"; 91 case Type::DoubleTyID: return ".f64"; 92 case Type::IntegerTyID: 93 switch (type->getPrimitiveSizeInBits()) { 94 default: llvm_unreachable("Unknown integer bit-width"); 95 case 16: return ".u16"; 96 case 32: return ".u32"; 97 case 64: return ".u64"; 98 } 99 case Type::ArrayTyID: 100 case Type::PointerTyID: 101 type = dyn_cast<SequentialType>(type)->getElementType(); 102 break; 103 } 104 } 105 return NULL; 106 } 107 108 bool PTXAsmPrinter::doFinalization(Module &M) { 109 // XXX Temproarily remove global variables so that doFinalization() will not 110 // emit them again (global variables are emitted at beginning). 111 112 Module::GlobalListType &global_list = M.getGlobalList(); 113 int i, n = global_list.size(); 114 GlobalVariable **gv_array = new GlobalVariable* [n]; 115 116 // first, back-up GlobalVariable in gv_array 117 i = 0; 118 for (Module::global_iterator I = global_list.begin(), E = global_list.end(); 119 I != E; ++I) 120 gv_array[i++] = &*I; 121 122 // second, empty global_list 123 while (!global_list.empty()) 124 global_list.remove(global_list.begin()); 125 126 // call doFinalization 127 bool ret = AsmPrinter::doFinalization(M); 128 129 // now we restore global variables 130 for (i = 0; i < n; i ++) 131 global_list.insert(global_list.end(), gv_array[i]); 132 133 delete[] gv_array; 134 return ret; 135 } 136 137 void PTXAsmPrinter::EmitStartOfAsmFile(Module &M) 138 { 139 const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>(); 140 141 // Emit the PTX .version and .target attributes 142 OutStreamer.EmitRawText(Twine("\t.version " + ST.getPTXVersionString())); 143 OutStreamer.EmitRawText(Twine("\t.target " + ST.getTargetString() + 144 (ST.supportsDouble() ? "" 145 : ", map_f64_to_f32"))); 146 // .address_size directive is optional, but it must immediately follow 147 // the .target directive if present within a module 148 if (ST.supportsPTX23()) { 149 std::string addrSize = ST.is64Bit() ? "64" : "32"; 150 OutStreamer.EmitRawText(Twine("\t.address_size " + addrSize)); 151 } 152 153 OutStreamer.AddBlankLine(); 154 155 // Define any .file directives 156 DebugInfoFinder DbgFinder; 157 DbgFinder.processModule(M); 158 159 for (DebugInfoFinder::iterator I = DbgFinder.compile_unit_begin(), 160 E = DbgFinder.compile_unit_end(); I != E; ++I) { 161 DICompileUnit DIUnit(*I); 162 StringRef FN = DIUnit.getFilename(); 163 StringRef Dir = DIUnit.getDirectory(); 164 GetOrCreateSourceID(FN, Dir); 165 } 166 167 OutStreamer.AddBlankLine(); 168 169 // declare global variables 170 for (Module::const_global_iterator i = M.global_begin(), e = M.global_end(); 171 i != e; ++i) 172 EmitVariableDeclaration(i); 173 } 174 175 void PTXAsmPrinter::EmitFunctionBodyStart() { 176 OutStreamer.EmitRawText(Twine("{")); 177 178 const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>(); 179 const PTXParamManager &PM = MFI->getParamManager(); 180 181 // Print register definitions 182 std::string regDefs; 183 unsigned numRegs; 184 185 // pred 186 numRegs = MFI->getNumRegistersForClass(PTX::RegPredRegisterClass); 187 if(numRegs > 0) { 188 regDefs += "\t.reg .pred %p<"; 189 regDefs += utostr(numRegs); 190 regDefs += ">;\n"; 191 } 192 193 // i16 194 numRegs = MFI->getNumRegistersForClass(PTX::RegI16RegisterClass); 195 if(numRegs > 0) { 196 regDefs += "\t.reg .b16 %rh<"; 197 regDefs += utostr(numRegs); 198 regDefs += ">;\n"; 199 } 200 201 // i32 202 numRegs = MFI->getNumRegistersForClass(PTX::RegI32RegisterClass); 203 if(numRegs > 0) { 204 regDefs += "\t.reg .b32 %r<"; 205 regDefs += utostr(numRegs); 206 regDefs += ">;\n"; 207 } 208 209 // i64 210 numRegs = MFI->getNumRegistersForClass(PTX::RegI64RegisterClass); 211 if(numRegs > 0) { 212 regDefs += "\t.reg .b64 %rd<"; 213 regDefs += utostr(numRegs); 214 regDefs += ">;\n"; 215 } 216 217 // f32 218 numRegs = MFI->getNumRegistersForClass(PTX::RegF32RegisterClass); 219 if(numRegs > 0) { 220 regDefs += "\t.reg .f32 %f<"; 221 regDefs += utostr(numRegs); 222 regDefs += ">;\n"; 223 } 224 225 // f64 226 numRegs = MFI->getNumRegistersForClass(PTX::RegF64RegisterClass); 227 if(numRegs > 0) { 228 regDefs += "\t.reg .f64 %fd<"; 229 regDefs += utostr(numRegs); 230 regDefs += ">;\n"; 231 } 232 233 // Local params 234 for (PTXParamManager::param_iterator i = PM.local_begin(), e = PM.local_end(); 235 i != e; ++i) { 236 regDefs += "\t.param .b"; 237 regDefs += utostr(PM.getParamSize(*i)); 238 regDefs += " "; 239 regDefs += PM.getParamName(*i); 240 regDefs += ";\n"; 241 } 242 243 OutStreamer.EmitRawText(Twine(regDefs)); 244 245 246 const MachineFrameInfo* FrameInfo = MF->getFrameInfo(); 247 DEBUG(dbgs() << "Have " << FrameInfo->getNumObjects() 248 << " frame object(s)\n"); 249 for (unsigned i = 0, e = FrameInfo->getNumObjects(); i != e; ++i) { 250 DEBUG(dbgs() << "Size of object: " << FrameInfo->getObjectSize(i) << "\n"); 251 if (FrameInfo->getObjectSize(i) > 0) { 252 std::string def = "\t.local .align "; 253 def += utostr(FrameInfo->getObjectAlignment(i)); 254 def += " .b8"; 255 def += " __local"; 256 def += utostr(i); 257 def += "["; 258 def += utostr(FrameInfo->getObjectSize(i)); // Convert to bits 259 def += "]"; 260 def += ";"; 261 OutStreamer.EmitRawText(Twine(def)); 262 } 263 } 264 265 //unsigned Index = 1; 266 // Print parameter passing params 267 //for (PTXMachineFunctionInfo::param_iterator 268 // i = MFI->paramBegin(), e = MFI->paramEnd(); i != e; ++i) { 269 // std::string def = "\t.param .b"; 270 // def += utostr(*i); 271 // def += " __ret_"; 272 // def += utostr(Index); 273 // Index++; 274 // def += ";"; 275 // OutStreamer.EmitRawText(Twine(def)); 276 //} 277 } 278 279 void PTXAsmPrinter::EmitFunctionBodyEnd() { 280 OutStreamer.EmitRawText(Twine("}")); 281 } 282 283 void PTXAsmPrinter::EmitInstruction(const MachineInstr *MI) { 284 MCInst TmpInst; 285 LowerPTXMachineInstrToMCInst(MI, TmpInst, *this); 286 OutStreamer.EmitInstruction(TmpInst); 287 } 288 289 void PTXAsmPrinter::EmitVariableDeclaration(const GlobalVariable *gv) { 290 // Check to see if this is a special global used by LLVM, if so, emit it. 291 if (EmitSpecialLLVMGlobal(gv)) 292 return; 293 294 MCSymbol *gvsym = Mang->getSymbol(gv); 295 296 assert(gvsym->isUndefined() && "Cannot define a symbol twice!"); 297 298 std::string decl; 299 300 // check if it is defined in some other translation unit 301 if (gv->isDeclaration()) 302 decl += ".extern "; 303 304 // state space: e.g., .global 305 decl += "."; 306 decl += getStateSpaceName(gv->getType()->getAddressSpace()); 307 decl += " "; 308 309 // alignment (optional) 310 unsigned alignment = gv->getAlignment(); 311 if (alignment != 0) { 312 decl += ".align "; 313 decl += utostr(gv->getAlignment()); 314 decl += " "; 315 } 316 317 318 if (PointerType::classof(gv->getType())) { 319 PointerType* pointerTy = dyn_cast<PointerType>(gv->getType()); 320 Type* elementTy = pointerTy->getElementType(); 321 322 decl += ".b8 "; 323 decl += gvsym->getName(); 324 decl += "["; 325 326 if (elementTy->isArrayTy()) 327 { 328 assert(elementTy->isArrayTy() && "Only pointers to arrays are supported"); 329 330 ArrayType* arrayTy = dyn_cast<ArrayType>(elementTy); 331 elementTy = arrayTy->getElementType(); 332 333 unsigned numElements = arrayTy->getNumElements(); 334 335 while (elementTy->isArrayTy()) { 336 337 arrayTy = dyn_cast<ArrayType>(elementTy); 338 elementTy = arrayTy->getElementType(); 339 340 numElements *= arrayTy->getNumElements(); 341 } 342 343 // FIXME: isPrimitiveType() == false for i16? 344 assert(elementTy->isSingleValueType() && 345 "Non-primitive types are not handled"); 346 347 // Compute the size of the array, in bytes. 348 uint64_t arraySize = (elementTy->getPrimitiveSizeInBits() >> 3) 349 * numElements; 350 351 decl += utostr(arraySize); 352 } 353 354 decl += "]"; 355 356 // handle string constants (assume ConstantArray means string) 357 358 if (gv->hasInitializer()) 359 { 360 const Constant *C = gv->getInitializer(); 361 if (const ConstantArray *CA = dyn_cast<ConstantArray>(C)) 362 { 363 decl += " = {"; 364 365 for (unsigned i = 0, e = C->getNumOperands(); i != e; ++i) 366 { 367 if (i > 0) decl += ","; 368 369 decl += "0x" + 370 utohexstr(cast<ConstantInt>(CA->getOperand(i))->getZExtValue()); 371 } 372 373 decl += "}"; 374 } 375 } 376 } 377 else { 378 // Note: this is currently the fall-through case and most likely generates 379 // incorrect code. 380 decl += getTypeName(gv->getType()); 381 decl += " "; 382 383 decl += gvsym->getName(); 384 385 if (ArrayType::classof(gv->getType()) || 386 PointerType::classof(gv->getType())) 387 decl += "[]"; 388 } 389 390 decl += ";"; 391 392 OutStreamer.EmitRawText(Twine(decl)); 393 394 OutStreamer.AddBlankLine(); 395 } 396 397 void PTXAsmPrinter::EmitFunctionEntryLabel() { 398 // The function label could have already been emitted if two symbols end up 399 // conflicting due to asm renaming. Detect this and emit an error. 400 if (!CurrentFnSym->isUndefined()) { 401 report_fatal_error("'" + Twine(CurrentFnSym->getName()) + 402 "' label emitted multiple times to assembly file"); 403 return; 404 } 405 406 const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>(); 407 const PTXParamManager &PM = MFI->getParamManager(); 408 const bool isKernel = MFI->isKernel(); 409 const PTXSubtarget& ST = TM.getSubtarget<PTXSubtarget>(); 410 const MachineRegisterInfo& MRI = MF->getRegInfo(); 411 412 std::string decl = isKernel ? ".entry" : ".func"; 413 414 if (!isKernel) { 415 decl += " ("; 416 if (ST.useParamSpaceForDeviceArgs()) { 417 for (PTXParamManager::param_iterator i = PM.ret_begin(), e = PM.ret_end(), 418 b = i; i != e; ++i) { 419 if (i != b) { 420 decl += ", "; 421 } 422 423 decl += ".param .b"; 424 decl += utostr(PM.getParamSize(*i)); 425 decl += " "; 426 decl += PM.getParamName(*i); 427 } 428 } else { 429 for (PTXMachineFunctionInfo::reg_iterator 430 i = MFI->retreg_begin(), e = MFI->retreg_end(), b = i; 431 i != e; ++i) { 432 if (i != b) { 433 decl += ", "; 434 } 435 decl += ".reg ."; 436 decl += getRegisterTypeName(*i, MRI); 437 decl += " "; 438 decl += MFI->getRegisterName(*i); 439 } 440 } 441 decl += ")"; 442 } 443 444 // Print function name 445 decl += " "; 446 decl += CurrentFnSym->getName().str(); 447 448 decl += " ("; 449 450 const Function *F = MF->getFunction(); 451 452 // Print parameters 453 if (isKernel || ST.useParamSpaceForDeviceArgs()) { 454 /*for (PTXParamManager::param_iterator i = PM.arg_begin(), e = PM.arg_end(), 455 b = i; i != e; ++i) { 456 if (i != b) { 457 decl += ", "; 458 } 459 460 decl += ".param .b"; 461 decl += utostr(PM.getParamSize(*i)); 462 decl += " "; 463 decl += PM.getParamName(*i); 464 }*/ 465 int Counter = 1; 466 for (Function::const_arg_iterator i = F->arg_begin(), e = F->arg_end(), 467 b = i; i != e; ++i) { 468 if (i != b) 469 decl += ", "; 470 const Type *ArgType = (*i).getType(); 471 decl += ".param .b"; 472 if (ArgType->isPointerTy()) { 473 if (ST.is64Bit()) 474 decl += "64"; 475 else 476 decl += "32"; 477 } else { 478 decl += utostr(ArgType->getPrimitiveSizeInBits()); 479 } 480 if (ArgType->isPointerTy() && ST.emitPtrAttribute()) { 481 const PointerType *PtrType = dyn_cast<const PointerType>(ArgType); 482 decl += " .ptr"; 483 switch (PtrType->getAddressSpace()) { 484 default: 485 llvm_unreachable("Unknown address space in argument"); 486 case PTXStateSpace::Global: 487 decl += " .global"; 488 break; 489 case PTXStateSpace::Shared: 490 decl += " .shared"; 491 break; 492 } 493 } 494 decl += " __param_"; 495 decl += utostr(Counter++); 496 } 497 } else { 498 for (PTXMachineFunctionInfo::reg_iterator 499 i = MFI->argreg_begin(), e = MFI->argreg_end(), b = i; 500 i != e; ++i) { 501 if (i != b) { 502 decl += ", "; 503 } 504 505 decl += ".reg ."; 506 decl += getRegisterTypeName(*i, MRI); 507 decl += " "; 508 decl += MFI->getRegisterName(*i); 509 } 510 } 511 decl += ")"; 512 513 OutStreamer.EmitRawText(Twine(decl)); 514 } 515 516 unsigned PTXAsmPrinter::GetOrCreateSourceID(StringRef FileName, 517 StringRef DirName) { 518 // If FE did not provide a file name, then assume stdin. 519 if (FileName.empty()) 520 return GetOrCreateSourceID("<stdin>", StringRef()); 521 522 // MCStream expects full path name as filename. 523 if (!DirName.empty() && !sys::path::is_absolute(FileName)) { 524 SmallString<128> FullPathName = DirName; 525 sys::path::append(FullPathName, FileName); 526 // Here FullPathName will be copied into StringMap by GetOrCreateSourceID. 527 return GetOrCreateSourceID(StringRef(FullPathName), StringRef()); 528 } 529 530 StringMapEntry<unsigned> &Entry = SourceIdMap.GetOrCreateValue(FileName); 531 if (Entry.getValue()) 532 return Entry.getValue(); 533 534 unsigned SrcId = SourceIdMap.size(); 535 Entry.setValue(SrcId); 536 537 // Print out a .file directive to specify files for .loc directives. 538 OutStreamer.EmitDwarfFileDirective(SrcId, "", Entry.getKey()); 539 540 return SrcId; 541 } 542 543 MCOperand PTXAsmPrinter::GetSymbolRef(const MachineOperand &MO, 544 const MCSymbol *Symbol) { 545 const MCExpr *Expr; 546 Expr = MCSymbolRefExpr::Create(Symbol, MCSymbolRefExpr::VK_None, OutContext); 547 return MCOperand::CreateExpr(Expr); 548 } 549 550 MCOperand PTXAsmPrinter::lowerOperand(const MachineOperand &MO) { 551 MCOperand MCOp; 552 const PTXMachineFunctionInfo *MFI = MF->getInfo<PTXMachineFunctionInfo>(); 553 const MCExpr *Expr; 554 const char *RegSymbolName; 555 switch (MO.getType()) { 556 default: 557 llvm_unreachable("Unknown operand type"); 558 case MachineOperand::MO_Register: 559 // We create register operands as symbols, since the PTXInstPrinter class 560 // has no way to map virtual registers back to a name without some ugly 561 // hacks. 562 // FIXME: Figure out a better way to handle virtual register naming. 563 RegSymbolName = MFI->getRegisterName(MO.getReg()); 564 Expr = MCSymbolRefExpr::Create(RegSymbolName, MCSymbolRefExpr::VK_None, 565 OutContext); 566 MCOp = MCOperand::CreateExpr(Expr); 567 break; 568 case MachineOperand::MO_Immediate: 569 MCOp = MCOperand::CreateImm(MO.getImm()); 570 break; 571 case MachineOperand::MO_MachineBasicBlock: 572 MCOp = MCOperand::CreateExpr(MCSymbolRefExpr::Create( 573 MO.getMBB()->getSymbol(), OutContext)); 574 break; 575 case MachineOperand::MO_GlobalAddress: 576 MCOp = GetSymbolRef(MO, Mang->getSymbol(MO.getGlobal())); 577 break; 578 case MachineOperand::MO_ExternalSymbol: 579 MCOp = GetSymbolRef(MO, GetExternalSymbolSymbol(MO.getSymbolName())); 580 break; 581 case MachineOperand::MO_FPImmediate: 582 APFloat Val = MO.getFPImm()->getValueAPF(); 583 bool ignored; 584 Val.convert(APFloat::IEEEdouble, APFloat::rmTowardZero, &ignored); 585 MCOp = MCOperand::CreateFPImm(Val.convertToDouble()); 586 break; 587 } 588 589 return MCOp; 590 } 591 592 // Force static initialization. 593 extern "C" void LLVMInitializePTXAsmPrinter() { 594 RegisterAsmPrinter<PTXAsmPrinter> X(ThePTX32Target); 595 RegisterAsmPrinter<PTXAsmPrinter> Y(ThePTX64Target); 596 } 597