From d9075301c381126c9f2ea1c07eaa7e48de29a3a2 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 9 Oct 2017 13:27:11 +0100 Subject: [PATCH 01/46] This adds cursory support for globals to the HIP module loading API. The style is purposefully alien so as to signal that HIP experts should turn it into HIP worthy code as soon as possible. [ROCm/clr commit: 582779c6c6cf4e34b0d72d13dad7ff1f9eaf205d] --- .../hip/hcc_detail/elfio/elf_types.hpp | 749 ++++++++++++++ .../include/hip/hcc_detail/elfio/elfio.hpp | 848 +++++++++++++++ .../hip/hcc_detail/elfio/elfio_dump.hpp | 972 ++++++++++++++++++ .../hip/hcc_detail/elfio/elfio_dynamic.hpp | 253 +++++ .../hip/hcc_detail/elfio/elfio_header.hpp | 146 +++ .../hip/hcc_detail/elfio/elfio_note.hpp | 166 +++ .../hip/hcc_detail/elfio/elfio_relocation.hpp | 369 +++++++ .../hip/hcc_detail/elfio/elfio_section.hpp | 296 ++++++ .../hip/hcc_detail/elfio/elfio_segment.hpp | 220 ++++ .../hip/hcc_detail/elfio/elfio_strings.hpp | 96 ++ .../hip/hcc_detail/elfio/elfio_symbols.hpp | 278 +++++ .../hip/hcc_detail/elfio/elfio_utils.hpp | 209 ++++ projects/clr/hipamd/src/hip_module.cpp | 306 ++++-- 13 files changed, 4804 insertions(+), 104 deletions(-) create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elf_types.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dump.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dynamic.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_header.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_note.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_relocation.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_section.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_segment.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_strings.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_symbols.hpp create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_utils.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elf_types.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elf_types.hpp new file mode 100644 index 0000000000..1b90c4c57d --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elf_types.hpp @@ -0,0 +1,749 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFTYPES_H +#define ELFTYPES_H + +#ifndef ELFIO_NO_OWN_TYPES + #if !defined(ELFIO_NO_CSTDINT) && !defined(ELFIO_NO_INTTYPES) + #include + #else + typedef unsigned char uint8_t; + typedef signed char int8_t; + typedef unsigned short uint16_t; + typedef signed short int16_t; + #ifdef _MSC_VER + typedef unsigned __int32 uint32_t; + typedef signed __int32 int32_t; + typedef unsigned __int64 uint64_t; + typedef signed __int64 int64_t; + #else + typedef unsigned int uint32_t; + typedef signed int int32_t; + typedef unsigned long long uint64_t; + typedef signed long long int64_t; + #endif // _MSC_VER + #endif // ELFIO_NO_CSTDINT +#endif // ELFIO_NO_OWN_TYPES + +namespace ELFIO { + +// Attention! Platform depended definitions. +typedef uint16_t Elf_Half; +typedef uint32_t Elf_Word; +typedef int32_t Elf_Sword; +typedef uint64_t Elf_Xword; +typedef int64_t Elf_Sxword; + +typedef uint32_t Elf32_Addr; +typedef uint32_t Elf32_Off; +typedef uint64_t Elf64_Addr; +typedef uint64_t Elf64_Off; + +#define Elf32_Half Elf_Half +#define Elf64_Half Elf_Half +#define Elf32_Word Elf_Word +#define Elf64_Word Elf_Word +#define Elf32_Sword Elf_Sword +#define Elf64_Sword Elf_Sword + +/////////////////////// +// ELF Header Constants + +// File type +#define ET_NONE 0 +#define ET_REL 1 +#define ET_EXEC 2 +#define ET_DYN 3 +#define ET_CORE 4 +#define ET_LOOS 0xFE00 +#define ET_HIOS 0xFEFF +#define ET_LOPROC 0xFF00 +#define ET_HIPROC 0xFFFF + + +#define EM_NONE 0 // No machine +#define EM_M32 1 // AT&T WE 32100 +#define EM_SPARC 2 // SUN SPARC +#define EM_386 3 // Intel 80386 +#define EM_68K 4 // Motorola m68k family +#define EM_88K 5 // Motorola m88k family +#define EM_486 6 // Intel 80486// Reserved for future use +#define EM_860 7 // Intel 80860 +#define EM_MIPS 8 // MIPS R3000 (officially, big-endian only) +#define EM_S370 9 // IBM System/370 +#define EM_MIPS_RS3_LE 10 // MIPS R3000 little-endian (Oct 4 1999 Draft) Deprecated +#define EM_res011 11 // Reserved +#define EM_res012 12 // Reserved +#define EM_res013 13 // Reserved +#define EM_res014 14 // Reserved +#define EM_PARISC 15 // HPPA +#define EM_res016 16 // Reserved +#define EM_VPP550 17 // Fujitsu VPP500 +#define EM_SPARC32PLUS 18 // Sun's "v8plus" +#define EM_960 19 // Intel 80960 +#define EM_PPC 20 // PowerPC +#define EM_PPC64 21 // 64-bit PowerPC +#define EM_S390 22 // IBM S/390 +#define EM_SPU 23 // Sony/Toshiba/IBM SPU +#define EM_res024 24 // Reserved +#define EM_res025 25 // Reserved +#define EM_res026 26 // Reserved +#define EM_res027 27 // Reserved +#define EM_res028 28 // Reserved +#define EM_res029 29 // Reserved +#define EM_res030 30 // Reserved +#define EM_res031 31 // Reserved +#define EM_res032 32 // Reserved +#define EM_res033 33 // Reserved +#define EM_res034 34 // Reserved +#define EM_res035 35 // Reserved +#define EM_V800 36 // NEC V800 series +#define EM_FR20 37 // Fujitsu FR20 +#define EM_RH32 38 // TRW RH32 +#define EM_MCORE 39 // Motorola M*Core // May also be taken by Fujitsu MMA +#define EM_RCE 39 // Old name for MCore +#define EM_ARM 40 // ARM +#define EM_OLD_ALPHA 41 // Digital Alpha +#define EM_SH 42 // Renesas (formerly Hitachi) / SuperH SH +#define EM_SPARCV9 43 // SPARC v9 64-bit +#define EM_TRICORE 44 // Siemens Tricore embedded processor +#define EM_ARC 45 // ARC Cores +#define EM_H8_300 46 // Renesas (formerly Hitachi) H8/300 +#define EM_H8_300H 47 // Renesas (formerly Hitachi) H8/300H +#define EM_H8S 48 // Renesas (formerly Hitachi) H8S +#define EM_H8_500 49 // Renesas (formerly Hitachi) H8/500 +#define EM_IA_64 50 // Intel IA-64 Processor +#define EM_MIPS_X 51 // Stanford MIPS-X +#define EM_COLDFIRE 52 // Motorola Coldfire +#define EM_68HC12 53 // Motorola M68HC12 +#define EM_MMA 54 // Fujitsu Multimedia Accelerator +#define EM_PCP 55 // Siemens PCP +#define EM_NCPU 56 // Sony nCPU embedded RISC processor +#define EM_NDR1 57 // Denso NDR1 microprocesspr +#define EM_STARCORE 58 // Motorola Star*Core processor +#define EM_ME16 59 // Toyota ME16 processor +#define EM_ST100 60 // STMicroelectronics ST100 processor +#define EM_TINYJ 61 // Advanced Logic Corp. TinyJ embedded processor +#define EM_X86_64 62 // Advanced Micro Devices X86-64 processor +#define EM_PDSP 63 // Sony DSP Processor +#define EM_PDP10 64 // Digital Equipment Corp. PDP-10 +#define EM_PDP11 65 // Digital Equipment Corp. PDP-11 +#define EM_FX66 66 // Siemens FX66 microcontroller +#define EM_ST9PLUS 67 // STMicroelectronics ST9+ 8/16 bit microcontroller +#define EM_ST7 68 // STMicroelectronics ST7 8-bit microcontroller +#define EM_68HC16 69 // Motorola MC68HC16 Microcontroller +#define EM_68HC11 70 // Motorola MC68HC11 Microcontroller +#define EM_68HC08 71 // Motorola MC68HC08 Microcontroller +#define EM_68HC05 72 // Motorola MC68HC05 Microcontroller +#define EM_SVX 73 // Silicon Graphics SVx +#define EM_ST19 74 // STMicroelectronics ST19 8-bit cpu +#define EM_VAX 75 // Digital VAX +#define EM_CRIS 76 // Axis Communications 32-bit embedded processor +#define EM_JAVELIN 77 // Infineon Technologies 32-bit embedded cpu +#define EM_FIREPATH 78 // Element 14 64-bit DSP processor +#define EM_ZSP 79 // LSI Logic's 16-bit DSP processor +#define EM_MMIX 80 // Donald Knuth's educational 64-bit processor +#define EM_HUANY 81 // Harvard's machine-independent format +#define EM_PRISM 82 // SiTera Prism +#define EM_AVR 83 // Atmel AVR 8-bit microcontroller +#define EM_FR30 84 // Fujitsu FR30 +#define EM_D10V 85 // Mitsubishi D10V +#define EM_D30V 86 // Mitsubishi D30V +#define EM_V850 87 // NEC v850 +#define EM_M32R 88 // Renesas M32R (formerly Mitsubishi M32R) +#define EM_MN10300 89 // Matsushita MN10300 +#define EM_MN10200 90 // Matsushita MN10200 +#define EM_PJ 91 // picoJava +#define EM_OPENRISC 92 // OpenRISC 32-bit embedded processor +#define EM_ARC_A5 93 // ARC Cores Tangent-A5 +#define EM_XTENSA 94 // Tensilica Xtensa Architecture +#define EM_VIDEOCORE 95 // Alphamosaic VideoCore processor +#define EM_TMM_GPP 96 // Thompson Multimedia General Purpose Processor +#define EM_NS32K 97 // National Semiconductor 32000 series +#define EM_TPC 98 // Tenor Network TPC processor +#define EM_SNP1K 99 // Trebia SNP 1000 processor +#define EM_ST200 100 // STMicroelectronics ST200 microcontroller +#define EM_IP2K 101 // Ubicom IP2022 micro controller +#define EM_MAX 102 // MAX Processor +#define EM_CR 103 // National Semiconductor CompactRISC +#define EM_F2MC16 104 // Fujitsu F2MC16 +#define EM_MSP430 105 // TI msp430 micro controller +#define EM_BLACKFIN 106 // ADI Blackfin +#define EM_SE_C33 107 // S1C33 Family of Seiko Epson processors +#define EM_SEP 108 // Sharp embedded microprocessor +#define EM_ARCA 109 // Arca RISC Microprocessor +#define EM_UNICORE 110 // Microprocessor series from PKU-Unity Ltd. and MPRC of Peking University +#define EM_EXCESS 111 // eXcess: 16/32/64-bit configurable embedded CPU +#define EM_DXP 112 // Icera Semiconductor Inc. Deep Execution Processor +#define EM_ALTERA_NIOS2 113 // Altera Nios II soft-core processor +#define EM_CRX 114 // National Semiconductor CRX +#define EM_XGATE 115 // Motorola XGATE embedded processor +#define EM_C166 116 // Infineon C16x/XC16x processor +#define EM_M16C 117 // Renesas M16C series microprocessors +#define EM_DSPIC30F 118 // Microchip Technology dsPIC30F Digital Signal Controller +#define EM_CE 119 // Freescale Communication Engine RISC core +#define EM_M32C 120 // Renesas M32C series microprocessors +#define EM_res121 121 // Reserved +#define EM_res122 122 // Reserved +#define EM_res123 123 // Reserved +#define EM_res124 124 // Reserved +#define EM_res125 125 // Reserved +#define EM_res126 126 // Reserved +#define EM_res127 127 // Reserved +#define EM_res128 128 // Reserved +#define EM_res129 129 // Reserved +#define EM_res130 130 // Reserved +#define EM_TSK3000 131 // Altium TSK3000 core +#define EM_RS08 132 // Freescale RS08 embedded processor +#define EM_res133 133 // Reserved +#define EM_ECOG2 134 // Cyan Technology eCOG2 microprocessor +#define EM_SCORE 135 // Sunplus Score +#define EM_SCORE7 135 // Sunplus S+core7 RISC processor +#define EM_DSP24 136 // New Japan Radio (NJR) 24-bit DSP Processor +#define EM_VIDEOCORE3 137 // Broadcom VideoCore III processor +#define EM_LATTICEMICO32 138 // RISC processor for Lattice FPGA architecture +#define EM_SE_C17 139 // Seiko Epson C17 family +#define EM_TI_C6000 140 // Texas Instruments TMS320C6000 DSP family +#define EM_TI_C2000 141 // Texas Instruments TMS320C2000 DSP family +#define EM_TI_C5500 142 // Texas Instruments TMS320C55x DSP family +#define EM_res143 143 // Reserved +#define EM_res144 144 // Reserved +#define EM_res145 145 // Reserved +#define EM_res146 146 // Reserved +#define EM_res147 147 // Reserved +#define EM_res148 148 // Reserved +#define EM_res149 149 // Reserved +#define EM_res150 150 // Reserved +#define EM_res151 151 // Reserved +#define EM_res152 152 // Reserved +#define EM_res153 153 // Reserved +#define EM_res154 154 // Reserved +#define EM_res155 155 // Reserved +#define EM_res156 156 // Reserved +#define EM_res157 157 // Reserved +#define EM_res158 158 // Reserved +#define EM_res159 159 // Reserved +#define EM_MMDSP_PLUS 160 // STMicroelectronics 64bit VLIW Data Signal Processor +#define EM_CYPRESS_M8C 161 // Cypress M8C microprocessor +#define EM_R32C 162 // Renesas R32C series microprocessors +#define EM_TRIMEDIA 163 // NXP Semiconductors TriMedia architecture family +#define EM_QDSP6 164 // QUALCOMM DSP6 Processor +#define EM_8051 165 // Intel 8051 and variants +#define EM_STXP7X 166 // STMicroelectronics STxP7x family +#define EM_NDS32 167 // Andes Technology compact code size embedded RISC processor family +#define EM_ECOG1 168 // Cyan Technology eCOG1X family +#define EM_ECOG1X 168 // Cyan Technology eCOG1X family +#define EM_MAXQ30 169 // Dallas Semiconductor MAXQ30 Core Micro-controllers +#define EM_XIMO16 170 // New Japan Radio (NJR) 16-bit DSP Processor +#define EM_MANIK 171 // M2000 Reconfigurable RISC Microprocessor +#define EM_CRAYNV2 172 // Cray Inc. NV2 vector architecture +#define EM_RX 173 // Renesas RX family +#define EM_METAG 174 // Imagination Technologies META processor architecture +#define EM_MCST_ELBRUS 175 // MCST Elbrus general purpose hardware architecture +#define EM_ECOG16 176 // Cyan Technology eCOG16 family +#define EM_CR16 177 // National Semiconductor CompactRISC 16-bit processor +#define EM_ETPU 178 // Freescale Extended Time Processing Unit +#define EM_SLE9X 179 // Infineon Technologies SLE9X core +#define EM_L1OM 180 // Intel L1OM +#define EM_INTEL181 181 // Reserved by Intel +#define EM_INTEL182 182 // Reserved by Intel +#define EM_res183 183 // Reserved by ARM +#define EM_res184 184 // Reserved by ARM +#define EM_AVR32 185 // Atmel Corporation 32-bit microprocessor family +#define EM_STM8 186 // STMicroeletronics STM8 8-bit microcontroller +#define EM_TILE64 187 // Tilera TILE64 multicore architecture family +#define EM_TILEPRO 188 // Tilera TILEPro multicore architecture family +#define EM_MICROBLAZE 189 // Xilinx MicroBlaze 32-bit RISC soft processor core +#define EM_CUDA 190 // NVIDIA CUDA architecture +#define EM_TILEGX 191 // Tilera TILE-Gx multicore architecture family +#define EM_CLOUDSHIELD 192 // CloudShield architecture family +#define EM_COREA_1ST 193 // KIPO-KAIST Core-A 1st generation processor family +#define EM_COREA_2ND 194 // KIPO-KAIST Core-A 2nd generation processor family +#define EM_ARC_COMPACT2 195 // Synopsys ARCompact V2 +#define EM_OPEN8 196 // Open8 8-bit RISC soft processor core +#define EM_RL78 197 // Renesas RL78 family +#define EM_VIDEOCORE5 198 // Broadcom VideoCore V processor +#define EM_78KOR 199 // Renesas 78KOR family +#define EM_56800EX 200 // Freescale 56800EX Digital Signal Controller (DSC) +#define EM_BA1 201 // Beyond BA1 CPU architecture +#define EM_BA2 202 // Beyond BA2 CPU architecture +#define EM_XCORE 203 // XMOS xCORE processor family +#define EM_MCHP_PIC 204 // Microchip 8-bit PIC(r) family +#define EM_INTEL205 205 // Reserved by Intel +#define EM_INTEL206 206 // Reserved by Intel +#define EM_INTEL207 207 // Reserved by Intel +#define EM_INTEL208 208 // Reserved by Intel +#define EM_INTEL209 209 // Reserved by Intel +#define EM_KM32 210 // KM211 KM32 32-bit processor +#define EM_KMX32 211 // KM211 KMX32 32-bit processor +#define EM_KMX16 212 // KM211 KMX16 16-bit processor +#define EM_KMX8 213 // KM211 KMX8 8-bit processor +#define EM_KVARC 214 // KM211 KVARC processor +#define EM_CDP 215 // Paneve CDP architecture family +#define EM_COGE 216 // Cognitive Smart Memory Processor +#define EM_COOL 217 // iCelero CoolEngine +#define EM_NORC 218 // Nanoradio Optimized RISC +#define EM_CSR_KALIMBA 219 // CSR Kalimba architecture family +#define EM_Z80 220 // Zilog Z80 +#define EM_VISIUM 221 // Controls and Data Services VISIUMcore processor +#define EM_FT32 222 // FTDI Chip FT32 high performance 32-bit RISC architecture +#define EM_MOXIE 223 // Moxie processor family +#define EM_AMDGPU 224 // AMD GPU architecture +#define EM_RISCV 243 // RISC-V +#define EM_LANAI 244 // Lanai processor +#define EM_CEVA 245 // CEVA Processor Architecture Family +#define EM_CEVA_X2 246 // CEVA X2 Processor Family +#define EM_BPF 247 // Linux BPF – in-kernel virtual machine + +// File version +#define EV_NONE 0 +#define EV_CURRENT 1 + +// Identification index +#define EI_MAG0 0 +#define EI_MAG1 1 +#define EI_MAG2 2 +#define EI_MAG3 3 +#define EI_CLASS 4 +#define EI_DATA 5 +#define EI_VERSION 6 +#define EI_OSABI 7 +#define EI_ABIVERSION 8 +#define EI_PAD 9 +#define EI_NIDENT 16 + +// Magic number +#define ELFMAG0 0x7F +#define ELFMAG1 'E' +#define ELFMAG2 'L' +#define ELFMAG3 'F' + +// File class +#define ELFCLASSNONE 0 +#define ELFCLASS32 1 +#define ELFCLASS64 2 + +// Encoding +#define ELFDATANONE 0 +#define ELFDATA2LSB 1 +#define ELFDATA2MSB 2 + +// OS extensions +#define ELFOSABI_NONE 0 // No extensions or unspecified +#define ELFOSABI_HPUX 1 // Hewlett-Packard HP-UX +#define ELFOSABI_NETBSD 2 // NetBSD +#define ELFOSABI_LINUX 3 // Linux +#define ELFOSABI_SOLARIS 6 // Sun Solaris +#define ELFOSABI_AIX 7 // AIX +#define ELFOSABI_IRIX 8 // IRIX +#define ELFOSABI_FREEBSD 9 // FreeBSD +#define ELFOSABI_TRU64 10 // Compaq TRU64 UNIX +#define ELFOSABI_MODESTO 11 // Novell Modesto +#define ELFOSABI_OPENBSD 12 // Open BSD +#define ELFOSABI_OPENVMS 13 // Open VMS +#define ELFOSABI_NSK 14 // Hewlett-Packard Non-Stop Kernel +#define ELFOSABI_AROS 15 // Amiga Research OS +#define ELFOSABI_FENIXOS 16 // The FenixOS highly scalable multi-core OS +// 64-255 Architecture-specific value range + + + +///////////////////// +// Sections constants + +// Section indexes +#define SHN_UNDEF 0 +#define SHN_LORESERVE 0xFF00 +#define SHN_LOPROC 0xFF00 +#define SHN_HIPROC 0xFF1F +#define SHN_LOOS 0xFF20 +#define SHN_HIOS 0xFF3F +#define SHN_ABS 0xFFF1 +#define SHN_COMMON 0xFFF2 +#define SHN_XINDEX 0xFFFF +#define SHN_HIRESERVE 0xFFFF + +// Section types +#define SHT_NULL 0 +#define SHT_PROGBITS 1 +#define SHT_SYMTAB 2 +#define SHT_STRTAB 3 +#define SHT_RELA 4 +#define SHT_HASH 5 +#define SHT_DYNAMIC 6 +#define SHT_NOTE 7 +#define SHT_NOBITS 8 +#define SHT_REL 9 +#define SHT_SHLIB 10 +#define SHT_DYNSYM 11 +#define SHT_INIT_ARRAY 14 +#define SHT_FINI_ARRAY 15 +#define SHT_PREINIT_ARRAY 16 +#define SHT_GROUP 17 +#define SHT_SYMTAB_SHNDX 18 +#define SHT_LOOS 0x60000000 +#define SHT_HIOS 0x6fffffff +#define SHT_LOPROC 0x70000000 +#define SHT_HIPROC 0x7FFFFFFF +#define SHT_LOUSER 0x80000000 +#define SHT_HIUSER 0xFFFFFFFF + +// Section attribute flags +#define SHF_WRITE 0x1 +#define SHF_ALLOC 0x2 +#define SHF_EXECINSTR 0x4 +#define SHF_MERGE 0x10 +#define SHF_STRINGS 0x20 +#define SHF_INFO_LINK 0x40 +#define SHF_LINK_ORDER 0x80 +#define SHF_OS_NONCONFORMING 0x100 +#define SHF_GROUP 0x200 +#define SHF_TLS 0x400 +#define SHF_MASKOS 0x0ff00000 +#define SHF_MASKPROC 0xF0000000 + +// Section group flags +#define GRP_COMDAT 0x1 +#define GRP_MASKOS 0x0ff00000 +#define GRP_MASKPROC 0xf0000000 + +// Symbol binding +#define STB_LOCAL 0 +#define STB_GLOBAL 1 +#define STB_WEAK 2 +#define STB_LOOS 10 +#define STB_HIOS 12 +#define STB_MULTIDEF 13 +#define STB_LOPROC 13 +#define STB_HIPROC 15 + +// Symbol types +#define STT_NOTYPE 0 +#define STT_OBJECT 1 +#define STT_FUNC 2 +#define STT_SECTION 3 +#define STT_FILE 4 +#define STT_COMMON 5 +#define STT_TLS 6 +#define STT_LOOS 10 +#define STT_HIOS 12 +#define STT_LOPROC 13 +#define STT_HIPROC 15 + +// Symbol visibility +#define STV_DEFAULT 0 +#define STV_INTERNAL 1 +#define STV_HIDDEN 2 +#define STV_PROTECTED 3 + +// Undefined name +#define STN_UNDEF 0 + +// Relocation types +#define R_386_NONE 0 +#define R_X86_64_NONE 0 +#define R_386_32 1 +#define R_X86_64_64 1 +#define R_386_PC32 2 +#define R_X86_64_PC32 2 +#define R_386_GOT32 3 +#define R_X86_64_GOT32 3 +#define R_386_PLT32 4 +#define R_X86_64_PLT32 4 +#define R_386_COPY 5 +#define R_X86_64_COPY 5 +#define R_386_GLOB_DAT 6 +#define R_X86_64_GLOB_DAT 6 +#define R_386_JMP_SLOT 7 +#define R_X86_64_JUMP_SLOT 7 +#define R_386_RELATIVE 8 +#define R_X86_64_RELATIVE 8 +#define R_386_GOTOFF 9 +#define R_X86_64_GOTPCREL 9 +#define R_386_GOTPC 10 +#define R_X86_64_32 10 +#define R_X86_64_32S 11 +#define R_X86_64_16 12 +#define R_X86_64_PC16 13 +#define R_X86_64_8 14 +#define R_X86_64_PC8 15 +#define R_X86_64_DTPMOD64 16 +#define R_X86_64_DTPOFF64 17 +#define R_X86_64_TPOFF64 18 +#define R_X86_64_TLSGD 19 +#define R_X86_64_TLSLD 20 +#define R_X86_64_DTPOFF32 21 +#define R_X86_64_GOTTPOFF 22 +#define R_X86_64_TPOFF32 23 +#define R_X86_64_PC64 24 +#define R_X86_64_GOTOFF64 25 +#define R_X86_64_GOTPC32 26 +#define R_X86_64_GOT64 27 +#define R_X86_64_GOTPCREL64 28 +#define R_X86_64_GOTPC64 29 +#define R_X86_64_GOTPLT64 30 +#define R_X86_64_PLTOFF64 31 +#define R_X86_64_GOTPC32_TLSDESC 34 +#define R_X86_64_TLSDESC_CALL 35 +#define R_X86_64_TLSDESC 36 +#define R_X86_64_IRELATIVE 37 +#define R_X86_64_GNU_VTINHERIT 250 +#define R_X86_64_GNU_VTENTRY 251 + +// Segment types +#define PT_NULL 0 +#define PT_LOAD 1 +#define PT_DYNAMIC 2 +#define PT_INTERP 3 +#define PT_NOTE 4 +#define PT_SHLIB 5 +#define PT_PHDR 6 +#define PT_TLS 7 +#define PT_LOOS 0x60000000 +#define PT_HIOS 0x6fffffff +#define PT_LOPROC 0x70000000 +#define PT_HIPROC 0x7FFFFFFF + +// Segment flags +#define PF_X 1 // Execute +#define PF_W 2 // Write +#define PF_R 4 // Read +#define PF_MASKOS 0x0ff00000 // Unspecified +#define PF_MASKPROC 0xf0000000 // Unspecified + +// Dynamic Array Tags +#define DT_NULL 0 +#define DT_NEEDED 1 +#define DT_PLTRELSZ 2 +#define DT_PLTGOT 3 +#define DT_HASH 4 +#define DT_STRTAB 5 +#define DT_SYMTAB 6 +#define DT_RELA 7 +#define DT_RELASZ 8 +#define DT_RELAENT 9 +#define DT_STRSZ 10 +#define DT_SYMENT 11 +#define DT_INIT 12 +#define DT_FINI 13 +#define DT_SONAME 14 +#define DT_RPATH 15 +#define DT_SYMBOLIC 16 +#define DT_REL 17 +#define DT_RELSZ 18 +#define DT_RELENT 19 +#define DT_PLTREL 20 +#define DT_DEBUG 21 +#define DT_TEXTREL 22 +#define DT_JMPREL 23 +#define DT_BIND_NOW 24 +#define DT_INIT_ARRAY 25 +#define DT_FINI_ARRAY 26 +#define DT_INIT_ARRAYSZ 27 +#define DT_FINI_ARRAYSZ 28 +#define DT_RUNPATH 29 +#define DT_FLAGS 30 +#define DT_ENCODING 32 +#define DT_PREINIT_ARRAY 32 +#define DT_PREINIT_ARRAYSZ 33 +#define DT_MAXPOSTAGS 34 +#define DT_LOOS 0x6000000D +#define DT_HIOS 0x6ffff000 +#define DT_LOPROC 0x70000000 +#define DT_HIPROC 0x7FFFFFFF + +// DT_FLAGS values +#define DF_ORIGIN 0x1 +#define DF_SYMBOLIC 0x2 +#define DF_TEXTREL 0x4 +#define DF_BIND_NOW 0x8 +#define DF_STATIC_TLS 0x10 + + +// ELF file header +struct Elf32_Ehdr { + unsigned char e_ident[EI_NIDENT]; + Elf_Half e_type; + Elf_Half e_machine; + Elf_Word e_version; + Elf32_Addr e_entry; + Elf32_Off e_phoff; + Elf32_Off e_shoff; + Elf_Word e_flags; + Elf_Half e_ehsize; + Elf_Half e_phentsize; + Elf_Half e_phnum; + Elf_Half e_shentsize; + Elf_Half e_shnum; + Elf_Half e_shstrndx; +}; + +struct Elf64_Ehdr { + unsigned char e_ident[EI_NIDENT]; + Elf_Half e_type; + Elf_Half e_machine; + Elf_Word e_version; + Elf64_Addr e_entry; + Elf64_Off e_phoff; + Elf64_Off e_shoff; + Elf_Word e_flags; + Elf_Half e_ehsize; + Elf_Half e_phentsize; + Elf_Half e_phnum; + Elf_Half e_shentsize; + Elf_Half e_shnum; + Elf_Half e_shstrndx; +}; + + +// Section header +struct Elf32_Shdr { + Elf_Word sh_name; + Elf_Word sh_type; + Elf_Word sh_flags; + Elf32_Addr sh_addr; + Elf32_Off sh_offset; + Elf_Word sh_size; + Elf_Word sh_link; + Elf_Word sh_info; + Elf_Word sh_addralign; + Elf_Word sh_entsize; +}; + +struct Elf64_Shdr { + Elf_Word sh_name; + Elf_Word sh_type; + Elf_Xword sh_flags; + Elf64_Addr sh_addr; + Elf64_Off sh_offset; + Elf_Xword sh_size; + Elf_Word sh_link; + Elf_Word sh_info; + Elf_Xword sh_addralign; + Elf_Xword sh_entsize; +}; + + +// Segment header +struct Elf32_Phdr { + Elf_Word p_type; + Elf32_Off p_offset; + Elf32_Addr p_vaddr; + Elf32_Addr p_paddr; + Elf_Word p_filesz; + Elf_Word p_memsz; + Elf_Word p_flags; + Elf_Word p_align; +}; + +struct Elf64_Phdr { + Elf_Word p_type; + Elf_Word p_flags; + Elf64_Off p_offset; + Elf64_Addr p_vaddr; + Elf64_Addr p_paddr; + Elf_Xword p_filesz; + Elf_Xword p_memsz; + Elf_Xword p_align; +}; + + +// Symbol table entry +struct Elf32_Sym { + Elf_Word st_name; + Elf32_Addr st_value; + Elf_Word st_size; + unsigned char st_info; + unsigned char st_other; + Elf_Half st_shndx; +}; + +struct Elf64_Sym { + Elf_Word st_name; + unsigned char st_info; + unsigned char st_other; + Elf_Half st_shndx; + Elf64_Addr st_value; + Elf_Xword st_size; +}; + + +#define ELF_ST_BIND(i) ((i)>>4) +#define ELF_ST_TYPE(i) ((i)&0xf) +#define ELF_ST_INFO(b,t) (((b)<<4)+((t)&0xf)) + +#define ELF_ST_VISIBILITY(o) ((o)&0x3) + + +// Relocation entries +struct Elf32_Rel { + Elf32_Addr r_offset; + Elf_Word r_info; +}; + +struct Elf32_Rela { + Elf32_Addr r_offset; + Elf_Word r_info; + Elf_Sword r_addend; +}; + +struct Elf64_Rel { + Elf64_Addr r_offset; + Elf_Xword r_info; +}; + +struct Elf64_Rela { + Elf64_Addr r_offset; + Elf_Xword r_info; + Elf_Sxword r_addend; +}; + + +#define ELF32_R_SYM(i) ((i)>>8) +#define ELF32_R_TYPE(i) ((unsigned char)(i)) +#define ELF32_R_INFO(s,t) (((s)<<8 )+(unsigned char)(t)) + +#define ELF64_R_SYM(i) ((i)>>32) +#define ELF64_R_TYPE(i) ((i)&0xffffffffL) +#define ELF64_R_INFO(s,t) ((((int64_t)s)<<32)+((t)&0xffffffffL)) + +// Dynamic structure +struct Elf32_Dyn { + Elf_Sword d_tag; + union { + Elf_Word d_val; + Elf32_Addr d_ptr; + } d_un; +}; + +struct Elf64_Dyn { + Elf_Sxword d_tag; + union { + Elf_Xword d_val; + Elf64_Addr d_ptr; + } d_un; +}; + +} // namespace ELFIO + +#endif // ELFTYPES_H diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio.hpp new file mode 100644 index 0000000000..46ce570406 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio.hpp @@ -0,0 +1,848 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_HPP +#define ELFIO_HPP + +#ifdef _MSC_VER +#pragma warning ( push ) +#pragma warning(disable:4996) +#pragma warning(disable:4355) +#pragma warning(disable:4244) +#endif + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "elf_types.hpp" +#include "elfio_utils.hpp" +#include "elfio_header.hpp" +#include "elfio_section.hpp" +#include "elfio_segment.hpp" +#include "elfio_strings.hpp" + +#define ELFIO_HEADER_ACCESS_GET( TYPE, FNAME ) \ +TYPE \ +get_##FNAME() const \ +{ \ + return header->get_##FNAME(); \ +} + +#define ELFIO_HEADER_ACCESS_GET_SET( TYPE, FNAME ) \ +TYPE \ +get_##FNAME() const \ +{ \ + return header->get_##FNAME(); \ +} \ +void \ +set_##FNAME( TYPE val ) \ +{ \ + header->set_##FNAME( val ); \ +} \ + +namespace ELFIO { + +//------------------------------------------------------------------------------ +class elfio +{ + public: +//------------------------------------------------------------------------------ + elfio() : sections( this ), segments( this ) + { + header = 0; + current_file_pos = 0; + create( ELFCLASS32, ELFDATA2LSB ); + } + +//------------------------------------------------------------------------------ + ~elfio() + { + clean(); + } + +//------------------------------------------------------------------------------ + void create( unsigned char file_class, unsigned char encoding ) + { + clean(); + convertor.setup( encoding ); + header = create_header( file_class, encoding ); + create_mandatory_sections(); + } + +//------------------------------------------------------------------------------ + bool load( const std::string& file_name ) + { + std::ifstream stream; + stream.open( file_name.c_str(), std::ios::in | std::ios::binary ); + if ( !stream ) { + return false; + } + + return load(stream); + } + +//------------------------------------------------------------------------------ + bool load( std::istream &stream ) + { + clean(); + + unsigned char e_ident[EI_NIDENT]; + + // Read ELF file signature + stream.seekg( 0 ); + stream.read( reinterpret_cast( &e_ident ), sizeof( e_ident ) ); + + // Is it ELF file? + if ( stream.gcount() != sizeof( e_ident ) || + e_ident[EI_MAG0] != ELFMAG0 || + e_ident[EI_MAG1] != ELFMAG1 || + e_ident[EI_MAG2] != ELFMAG2 || + e_ident[EI_MAG3] != ELFMAG3 ) { + return false; + } + + if ( ( e_ident[EI_CLASS] != ELFCLASS64 ) && + ( e_ident[EI_CLASS] != ELFCLASS32 )) { + return false; + } + + convertor.setup( e_ident[EI_DATA] ); + + header = create_header( e_ident[EI_CLASS], e_ident[EI_DATA] ); + if ( 0 == header ) { + return false; + } + if ( !header->load( stream ) ) { + return false; + } + + load_sections( stream ); + load_segments( stream ); + + return true; + } + +//------------------------------------------------------------------------------ + bool save( const std::string& file_name ) + { + std::ofstream f( file_name.c_str(), std::ios::out | std::ios::binary ); + + if ( !f ) { + return false; + } + + bool is_still_good = true; + + // Define layout specific header fields + // The position of the segment table is fixed after the header. + // The position of the section table is variable and needs to be fixed + // before saving. + header->set_segments_num( segments.size() ); + header->set_segments_offset( segments.size() ? header->get_header_size() : 0 ); + header->set_sections_num( sections.size() ); + header->set_sections_offset( 0 ); + + // Layout the first section right after the segment table + current_file_pos = header->get_header_size() + + header->get_segment_entry_size() * header->get_segments_num(); + + is_still_good = layout_segments_and_their_sections(); + is_still_good = is_still_good && layout_sections_without_segments(); + is_still_good = is_still_good && layout_section_table(); + + is_still_good = is_still_good && save_header( f ); + is_still_good = is_still_good && save_sections( f ); + is_still_good = is_still_good && save_segments( f ); + + f.close(); + + return is_still_good; + } + +//------------------------------------------------------------------------------ + // ELF header access functions + ELFIO_HEADER_ACCESS_GET( unsigned char, class ); + ELFIO_HEADER_ACCESS_GET( unsigned char, elf_version ); + ELFIO_HEADER_ACCESS_GET( unsigned char, encoding ); + ELFIO_HEADER_ACCESS_GET( Elf_Word, version ); + ELFIO_HEADER_ACCESS_GET( Elf_Half, header_size ); + ELFIO_HEADER_ACCESS_GET( Elf_Half, section_entry_size ); + ELFIO_HEADER_ACCESS_GET( Elf_Half, segment_entry_size ); + + ELFIO_HEADER_ACCESS_GET_SET( unsigned char, os_abi ); + ELFIO_HEADER_ACCESS_GET_SET( unsigned char, abi_version ); + ELFIO_HEADER_ACCESS_GET_SET( Elf_Half, type ); + ELFIO_HEADER_ACCESS_GET_SET( Elf_Half, machine ); + ELFIO_HEADER_ACCESS_GET_SET( Elf_Word, flags ); + ELFIO_HEADER_ACCESS_GET_SET( Elf64_Addr, entry ); + ELFIO_HEADER_ACCESS_GET_SET( Elf64_Off, sections_offset ); + ELFIO_HEADER_ACCESS_GET_SET( Elf64_Off, segments_offset ); + ELFIO_HEADER_ACCESS_GET_SET( Elf_Half, section_name_str_index ); + +//------------------------------------------------------------------------------ + const endianess_convertor& get_convertor() const + { + return convertor; + } + +//------------------------------------------------------------------------------ + Elf_Xword get_default_entry_size( Elf_Word section_type ) const + { + switch( section_type ) { + case SHT_RELA: + if ( header->get_class() == ELFCLASS64 ) { + return sizeof( Elf64_Rela ); + } + else { + return sizeof( Elf32_Rela ); + } + case SHT_REL: + if ( header->get_class() == ELFCLASS64 ) { + return sizeof( Elf64_Rel ); + } + else { + return sizeof( Elf32_Rel ); + } + case SHT_SYMTAB: + if ( header->get_class() == ELFCLASS64 ) { + return sizeof( Elf64_Sym ); + } + else { + return sizeof( Elf32_Sym ); + } + case SHT_DYNAMIC: + if ( header->get_class() == ELFCLASS64 ) { + return sizeof( Elf64_Dyn ); + } + else { + return sizeof( Elf32_Dyn ); + } + default: + return 0; + } + } + +//------------------------------------------------------------------------------ + private: +//------------------------------------------------------------------------------ + void clean() + { + delete header; + header = 0; + + std::vector::const_iterator it; + for ( it = sections_.begin(); it != sections_.end(); ++it ) { + delete *it; + } + sections_.clear(); + + std::vector::const_iterator it1; + for ( it1 = segments_.begin(); it1 != segments_.end(); ++it1 ) { + delete *it1; + } + segments_.clear(); + } + +//------------------------------------------------------------------------------ + elf_header* create_header( unsigned char file_class, unsigned char encoding ) + { + elf_header* new_header = 0; + + if ( file_class == ELFCLASS64 ) { + new_header = new elf_header_impl< Elf64_Ehdr >( &convertor, + encoding ); + } + else if ( file_class == ELFCLASS32 ) { + new_header = new elf_header_impl< Elf32_Ehdr >( &convertor, + encoding ); + } + else { + return 0; + } + + return new_header; + } + +//------------------------------------------------------------------------------ + section* create_section() + { + section* new_section; + unsigned char file_class = get_class(); + + if ( file_class == ELFCLASS64 ) { + new_section = new section_impl( &convertor ); + } + else if ( file_class == ELFCLASS32 ) { + new_section = new section_impl( &convertor ); + } + else { + return 0; + } + + new_section->set_index( (Elf_Half)sections_.size() ); + sections_.push_back( new_section ); + + return new_section; + } + + +//------------------------------------------------------------------------------ + segment* create_segment() + { + segment* new_segment; + unsigned char file_class = header->get_class(); + + if ( file_class == ELFCLASS64 ) { + new_segment = new segment_impl( &convertor ); + } + else if ( file_class == ELFCLASS32 ) { + new_segment = new segment_impl( &convertor ); + } + else { + return 0; + } + + new_segment->set_index( (Elf_Half)segments_.size() ); + segments_.push_back( new_segment ); + + return new_segment; + } + +//------------------------------------------------------------------------------ + void create_mandatory_sections() + { + // Create null section without calling to 'add_section' as no string + // section containing section names exists yet + section* sec0 = create_section(); + sec0->set_index( 0 ); + sec0->set_name( "" ); + sec0->set_name_string_offset( 0 ); + + set_section_name_str_index( 1 ); + section* shstrtab = sections.add( ".shstrtab" ); + shstrtab->set_type( SHT_STRTAB ); + shstrtab->set_addr_align( 1 ); + } + +//------------------------------------------------------------------------------ + Elf_Half load_sections( std::istream& stream ) + { + Elf_Half entry_size = header->get_section_entry_size(); + Elf_Half num = header->get_sections_num(); + Elf64_Off offset = header->get_sections_offset(); + + for ( Elf_Half i = 0; i < num; ++i ) { + section* sec = create_section(); + sec->load( stream, (std::streamoff)offset + i * entry_size ); + sec->set_index( i ); + // To mark that the section is not permitted to reassign address + // during layout calculation + sec->set_address( sec->get_address() ); + } + + Elf_Half shstrndx = get_section_name_str_index(); + + if ( SHN_UNDEF != shstrndx ) { + string_section_accessor str_reader( sections[shstrndx] ); + for ( Elf_Half i = 0; i < num; ++i ) { + Elf_Word offset = sections[i]->get_name_string_offset(); + const char* p = str_reader.get_string( offset ); + if ( p != 0 ) { + sections[i]->set_name( p ); + } + } + } + + return num; + } + +//------------------------------------------------------------------------------ + bool load_segments( std::istream& stream ) + { + Elf_Half entry_size = header->get_segment_entry_size(); + Elf_Half num = header->get_segments_num(); + Elf64_Off offset = header->get_segments_offset(); + + for ( Elf_Half i = 0; i < num; ++i ) { + segment* seg; + unsigned char file_class = header->get_class(); + + if ( file_class == ELFCLASS64 ) { + seg = new segment_impl( &convertor ); + } + else if ( file_class == ELFCLASS32 ) { + seg = new segment_impl( &convertor ); + } + else { + return false; + } + + seg->load( stream, (std::streamoff)offset + i * entry_size ); + seg->set_index( i ); + + // Add sections to the segments (similar to readelfs algorithm) + Elf64_Off segBaseOffset = seg->get_offset(); + Elf64_Off segEndOffset = segBaseOffset + seg->get_file_size(); + Elf64_Off segVBaseAddr = seg->get_virtual_address(); + Elf64_Off segVEndAddr = segVBaseAddr + seg->get_memory_size(); + for( Elf_Half j = 0; j < sections.size(); ++j ) { + const section* psec = sections[j]; + + // SHF_ALLOC sections are matched based on the virtual address + // otherwise the file offset is matched + if( psec->get_flags() & SHF_ALLOC + ? (segVBaseAddr <= psec->get_address() + && psec->get_address() + psec->get_size() + <= segVEndAddr) + : (segBaseOffset <= psec->get_offset() + && psec->get_offset() + psec->get_size() + <= segEndOffset)) { + seg->add_section_index( psec->get_index(), + psec->get_addr_align() ); + } + } + + // Add section into the segments' container + segments_.push_back( seg ); + } + + return true; + } + +//------------------------------------------------------------------------------ + bool save_header( std::ofstream& f ) + { + return header->save( f ); + } + +//------------------------------------------------------------------------------ + bool save_sections( std::ofstream& f ) + { + for ( unsigned int i = 0; i < sections_.size(); ++i ) { + section *sec = sections_.at(i); + + std::streampos headerPosition = + (std::streamoff)header->get_sections_offset() + + header->get_section_entry_size() * sec->get_index(); + + sec->save(f,headerPosition,sec->get_offset()); + } + return true; + } + +//------------------------------------------------------------------------------ + bool save_segments( std::ofstream& f ) + { + for ( unsigned int i = 0; i < segments_.size(); ++i ) { + segment *seg = segments_.at(i); + + std::streampos headerPosition = header->get_segments_offset() + + header->get_segment_entry_size()*seg->get_index(); + + seg->save( f, headerPosition, seg->get_offset() ); + } + return true; + } + +//------------------------------------------------------------------------------ + bool is_section_without_segment( unsigned int section_index ) + { + bool found = false; + + for ( unsigned int j = 0; !found && ( j < segments.size() ); ++j ) { + for ( unsigned int k = 0; + !found && ( k < segments[j]->get_sections_num() ); + ++k ) { + found = segments[j]->get_section_index_at( k ) == section_index; + } + } + + return !found; + } + +//------------------------------------------------------------------------------ + bool is_subsequence_of( segment* seg1, segment* seg2 ) + { + // Return 'true' if sections of seg1 are a subset of sections in seg2 + const std::vector& sections1 = seg1->get_sections(); + const std::vector& sections2 = seg2->get_sections(); + + bool found = false; + if ( sections1.size() < sections2.size() ) { + found = std::includes( sections2.begin(), sections2.end(), + sections1.begin(), sections1.end() ); + } + + return found; + } + +//------------------------------------------------------------------------------ + std::vector get_ordered_segments( ) + { + std::vector res; + std::deque worklist; + + res.reserve(segments.size()); + std::copy( segments_.begin(), segments_.end(), + std::back_inserter( worklist )) ; + + // Bring the segments which start at address 0 to the front + size_t nextSlot = 0; + for( size_t i = 0; i < worklist.size(); ++i ) { + if( i != nextSlot && worklist[i]->is_offset_initialized() + && worklist[i]->get_offset() == 0 ) { + std::swap(worklist[i],worklist[nextSlot]); + ++nextSlot; + } + } + + while ( !worklist.empty() ) { + segment *seg = worklist.front(); + worklist.pop_front(); + + size_t i = 0; + for ( ; i < worklist.size(); ++i ) { + if ( is_subsequence_of( seg, worklist[i] ) ) { + break; + } + } + + if ( i < worklist.size() ) + worklist.push_back(seg); + else + res.push_back(seg); + } + + return res; + } + + +//------------------------------------------------------------------------------ + bool layout_sections_without_segments( ) + { + for ( unsigned int i = 0; i < sections_.size(); ++i ) { + if ( is_section_without_segment( i ) ) { + section *sec = sections_[i]; + + Elf_Xword section_align = sec->get_addr_align(); + if ( section_align > 1 && current_file_pos % section_align != 0 ) { + current_file_pos += section_align - + current_file_pos % section_align; + } + + if ( 0 != sec->get_index() ) + sec->set_offset(current_file_pos); + + if ( SHT_NOBITS != sec->get_type() && + SHT_NULL != sec->get_type() ) { + current_file_pos += sec->get_size(); + } + } + } + + return true; + } + + +//------------------------------------------------------------------------------ + bool layout_segments_and_their_sections( ) + { + std::vector worklist; + std::vector section_generated(sections.size(),false); + + // Get segments in a order in where segments which contain a + // sub sequence of other segments are located at the end + worklist = get_ordered_segments(); + + for ( unsigned int i = 0; i < worklist.size(); ++i ) { + Elf_Xword segment_memory = 0; + Elf_Xword segment_filesize = 0; + Elf_Xword seg_start_pos = current_file_pos; + segment* seg = worklist[i]; + + // Special case: PHDR segment + // This segment contains the program headers but no sections + if ( seg->get_type() == PT_PHDR && seg->get_sections_num() == 0 ) { + seg_start_pos = header->get_segments_offset(); + segment_memory = segment_filesize = + header->get_segment_entry_size() * header->get_segments_num(); + } + // Special case: + // Segments which start with the NULL section and have further sections + else if ( seg->get_sections_num() > 1 + && sections[seg->get_section_index_at( 0 )]->get_type() == SHT_NULL ) { + seg_start_pos = 0; + if ( seg->get_sections_num() ) { + segment_memory = segment_filesize = current_file_pos; + } + } + // New segments with not generated sections + // have to be aligned + else if ( seg->get_sections_num() + && !section_generated[seg->get_section_index_at( 0 )] ) { + Elf64_Off cur_page_alignment = current_file_pos % seg->get_align(); + Elf64_Off req_page_alignment = seg->get_virtual_address() % seg->get_align(); + Elf64_Off error = req_page_alignment - cur_page_alignment; + + current_file_pos += ( seg->get_align() + error ) % seg->get_align(); + seg_start_pos = current_file_pos; + } + else if ( seg->get_sections_num() ) { + seg_start_pos = sections[seg->get_section_index_at( 0 )]->get_offset(); + } + + // Write segment's data + for ( unsigned int j = 0; j < seg->get_sections_num(); ++j ) { + Elf_Half index = seg->get_section_index_at( j ); + + section* sec = sections[ index ]; + + // The NULL section is always generated + if ( SHT_NULL == sec->get_type()) { + section_generated[index] = true; + continue; + } + + Elf_Xword secAlign = 0; + // Fix up the alignment + if ( !section_generated[index] && sec->is_address_initialized() + && SHT_NOBITS != sec->get_type() + && SHT_NULL != sec->get_type() ) { + // Align the sections based on the virtual addresses + // when possible (this is what matters for execution) + Elf64_Off req_offset = sec->get_address() - seg->get_virtual_address(); + Elf64_Off cur_offset = current_file_pos - seg_start_pos; + secAlign = req_offset - cur_offset; + } + else if (!section_generated[index]) { + // If no address has been specified then only the section + // alignment constraint has to be matched + Elf_Xword align = sec->get_addr_align(); + if (align == 0) { + align = 1; + } + Elf64_Off error = current_file_pos % align; + secAlign = ( align - error ) % align; + } + else { + // Alignment for already generated sections + secAlign = sec->get_offset() - seg_start_pos - segment_filesize; + } + + // Determine the segment file and memory sizes + // Special case .tbss section (NOBITS) in non TLS segment + if ( (sec->get_flags() & SHF_ALLOC) + && !( (sec->get_flags() & SHF_TLS) && (seg->get_type() != PT_TLS) + && ( SHT_NOBITS == sec->get_type())) ) + segment_memory += sec->get_size() + secAlign; + if ( SHT_NOBITS != sec->get_type() && SHT_NULL != sec->get_type() ) + segment_filesize += sec->get_size() + secAlign; + + // Nothing to be done when generating nested segments + if(section_generated[index]) { + continue; + } + + current_file_pos += secAlign; + + // Set the section addresses when missing + if ( !sec->is_address_initialized() ) + sec->set_address( seg->get_virtual_address() + + current_file_pos - seg_start_pos); + + if ( 0 != sec->get_index() ) + sec->set_offset(current_file_pos); + + if ( SHT_NOBITS != sec->get_type() && SHT_NULL != sec->get_type() ) + current_file_pos += sec->get_size(); + section_generated[index] = true; + } + + seg->set_file_size( segment_filesize ); + seg->set_memory_size( segment_memory ); + seg->set_offset(seg_start_pos); + } + + return true; + } + +//------------------------------------------------------------------------------ + bool layout_section_table() + { + // Simply place the section table at the end for now + Elf64_Off alignmentError = current_file_pos % 4; + current_file_pos += ( 4 - alignmentError ) % 4; + header->set_sections_offset(current_file_pos); + return true; + } + + +//------------------------------------------------------------------------------ + public: + friend class Sections; + class Sections { + public: +//------------------------------------------------------------------------------ + Sections( elfio* parent_ ) : + parent( parent_ ) + { + } + +//------------------------------------------------------------------------------ + Elf_Half size() const + { + return (Elf_Half)parent->sections_.size(); + } + +//------------------------------------------------------------------------------ + section* operator[]( unsigned int index ) const + { + section* sec = 0; + + if ( index < parent->sections_.size() ) { + sec = parent->sections_[index]; + } + + return sec; + } + +//------------------------------------------------------------------------------ + section* operator[]( const std::string& name ) const + { + section* sec = 0; + + std::vector::const_iterator it; + for ( it = parent->sections_.begin(); + it != parent->sections_.end(); + ++it ) { + if ( (*it)->get_name() == name ) { + sec = *it; + break; + } + } + + return sec; + } + +//------------------------------------------------------------------------------ + section* add( const std::string& name ) + { + section* new_section = parent->create_section(); + new_section->set_name( name ); + + Elf_Half str_index = parent->get_section_name_str_index(); + section* string_table( parent->sections_[str_index] ); + string_section_accessor str_writer( string_table ); + Elf_Word pos = str_writer.add_string( name ); + new_section->set_name_string_offset( pos ); + + return new_section; + } + +//------------------------------------------------------------------------------ + std::vector::iterator begin() { + return parent->sections_.begin(); + } + +//------------------------------------------------------------------------------ + std::vector::iterator end() { + return parent->sections_.end(); + } + +//------------------------------------------------------------------------------ + private: + elfio* parent; + } sections; + +//------------------------------------------------------------------------------ + public: + friend class Segments; + class Segments { + public: +//------------------------------------------------------------------------------ + Segments( elfio* parent_ ) : + parent( parent_ ) + { + } + +//------------------------------------------------------------------------------ + Elf_Half size() const + { + return (Elf_Half)parent->segments_.size(); + } + +//------------------------------------------------------------------------------ + segment* operator[]( unsigned int index ) const + { + return parent->segments_[index]; + } + + +//------------------------------------------------------------------------------ + segment* add() + { + return parent->create_segment(); + } + +//------------------------------------------------------------------------------ + std::vector::iterator begin() { + return parent->segments_.begin(); + } + +//------------------------------------------------------------------------------ + std::vector::iterator end() { + return parent->segments_.end(); + } + +//------------------------------------------------------------------------------ + private: + elfio* parent; + } segments; + +//------------------------------------------------------------------------------ + private: + elf_header* header; + std::vector sections_; + std::vector segments_; + endianess_convertor convertor; + + Elf_Xword current_file_pos; +}; + +} // namespace ELFIO + +#include "elfio_symbols.hpp" +#include "elfio_note.hpp" +#include "elfio_relocation.hpp" +#include "elfio_dynamic.hpp" + +#ifdef _MSC_VER +#pragma warning ( pop ) +#endif + +#endif // ELFIO_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dump.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dump.hpp new file mode 100644 index 0000000000..0494852960 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dump.hpp @@ -0,0 +1,972 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_DUMP_HPP +#define ELFIO_DUMP_HPP + +#include +#include +#include +#include +#include +#include "elfio.hpp" + +namespace ELFIO { + + +static struct class_table_t { + const char key; + const char* str; +} class_table [] = +{ + { ELFCLASS32, "ELF32" }, + { ELFCLASS64, "ELF64" }, +}; + + +static struct endian_table_t { + const char key; + const char* str; +} endian_table [] = +{ + { ELFDATANONE, "None" }, + { ELFDATA2LSB, "Little endian" }, + { ELFDATA2MSB, "Big endian" }, +}; + + +static struct version_table_t { + const Elf64_Word key; + const char* str; +} version_table [] = +{ + { EV_NONE , "None" }, + { EV_CURRENT, "Current" }, +}; + + +static struct type_table_t { + const Elf32_Half key; + const char* str; +} type_table [] = +{ + { ET_NONE, "No file type" }, + { ET_REL , "Relocatable file" }, + { ET_EXEC, "Executable file" }, + { ET_DYN , "Shared object file" }, + { ET_CORE, "Core file" }, +}; + + +static struct machine_table_t { + const Elf64_Half key; + const char* str; +} machine_table [] = +{ + { EM_NONE , "No machine" }, + { EM_M32 , "AT&T WE 32100" }, + { EM_SPARC , "SUN SPARC" }, + { EM_386 , "Intel 80386" }, + { EM_68K , "Motorola m68k family" }, + { EM_88K , "Motorola m88k family" }, + { EM_486 , "Intel 80486// Reserved for future use" }, + { EM_860 , "Intel 80860" }, + { EM_MIPS , "MIPS R3000 (officially, big-endian only)" }, + { EM_S370 , "IBM System/370" }, + { EM_MIPS_RS3_LE , "MIPS R3000 little-endian (Oct 4 1999 Draft) Deprecated" }, + { EM_res011 , "Reserved" }, + { EM_res012 , "Reserved" }, + { EM_res013 , "Reserved" }, + { EM_res014 , "Reserved" }, + { EM_PARISC , "HPPA" }, + { EM_res016 , "Reserved" }, + { EM_VPP550 , "Fujitsu VPP500" }, + { EM_SPARC32PLUS , "Sun's v8plus" }, + { EM_960 , "Intel 80960" }, + { EM_PPC , "PowerPC" }, + { EM_PPC64 , "64-bit PowerPC" }, + { EM_S390 , "IBM S/390" }, + { EM_SPU , "Sony/Toshiba/IBM SPU" }, + { EM_res024 , "Reserved" }, + { EM_res025 , "Reserved" }, + { EM_res026 , "Reserved" }, + { EM_res027 , "Reserved" }, + { EM_res028 , "Reserved" }, + { EM_res029 , "Reserved" }, + { EM_res030 , "Reserved" }, + { EM_res031 , "Reserved" }, + { EM_res032 , "Reserved" }, + { EM_res033 , "Reserved" }, + { EM_res034 , "Reserved" }, + { EM_res035 , "Reserved" }, + { EM_V800 , "NEC V800 series" }, + { EM_FR20 , "Fujitsu FR20" }, + { EM_RH32 , "TRW RH32" }, + { EM_MCORE , "Motorola M*Core // May also be taken by Fujitsu MMA" }, + { EM_RCE , "Old name for MCore" }, + { EM_ARM , "ARM" }, + { EM_OLD_ALPHA , "Digital Alpha" }, + { EM_SH , "Renesas (formerly Hitachi) / SuperH SH" }, + { EM_SPARCV9 , "SPARC v9 64-bit" }, + { EM_TRICORE , "Siemens Tricore embedded processor" }, + { EM_ARC , "ARC Cores" }, + { EM_H8_300 , "Renesas (formerly Hitachi) H8/300" }, + { EM_H8_300H , "Renesas (formerly Hitachi) H8/300H" }, + { EM_H8S , "Renesas (formerly Hitachi) H8S" }, + { EM_H8_500 , "Renesas (formerly Hitachi) H8/500" }, + { EM_IA_64 , "Intel IA-64 Processor" }, + { EM_MIPS_X , "Stanford MIPS-X" }, + { EM_COLDFIRE , "Motorola Coldfire" }, + { EM_68HC12 , "Motorola M68HC12" }, + { EM_MMA , "Fujitsu Multimedia Accelerator" }, + { EM_PCP , "Siemens PCP" }, + { EM_NCPU , "Sony nCPU embedded RISC processor" }, + { EM_NDR1 , "Denso NDR1 microprocesspr" }, + { EM_STARCORE , "Motorola Star*Core processor" }, + { EM_ME16 , "Toyota ME16 processor" }, + { EM_ST100 , "STMicroelectronics ST100 processor" }, + { EM_TINYJ , "Advanced Logic Corp. TinyJ embedded processor" }, + { EM_X86_64 , "Advanced Micro Devices X86-64 processor" }, + { EM_PDSP , "Sony DSP Processor" }, + { EM_PDP10 , "Digital Equipment Corp. PDP-10" }, + { EM_PDP11 , "Digital Equipment Corp. PDP-11" }, + { EM_FX66 , "Siemens FX66 microcontroller" }, + { EM_ST9PLUS , "STMicroelectronics ST9+ 8/16 bit microcontroller" }, + { EM_ST7 , "STMicroelectronics ST7 8-bit microcontroller" }, + { EM_68HC16 , "Motorola MC68HC16 Microcontroller" }, + { EM_68HC11 , "Motorola MC68HC11 Microcontroller" }, + { EM_68HC08 , "Motorola MC68HC08 Microcontroller" }, + { EM_68HC05 , "Motorola MC68HC05 Microcontroller" }, + { EM_SVX , "Silicon Graphics SVx" }, + { EM_ST19 , "STMicroelectronics ST19 8-bit cpu" }, + { EM_VAX , "Digital VAX" }, + { EM_CRIS , "Axis Communications 32-bit embedded processor" }, + { EM_JAVELIN , "Infineon Technologies 32-bit embedded cpu" }, + { EM_FIREPATH , "Element 14 64-bit DSP processor" }, + { EM_ZSP , "LSI Logic's 16-bit DSP processor" }, + { EM_MMIX , "Donald Knuth's educational 64-bit processor" }, + { EM_HUANY , "Harvard's machine-independent format" }, + { EM_PRISM , "SiTera Prism" }, + { EM_AVR , "Atmel AVR 8-bit microcontroller" }, + { EM_FR30 , "Fujitsu FR30" }, + { EM_D10V , "Mitsubishi D10V" }, + { EM_D30V , "Mitsubishi D30V" }, + { EM_V850 , "NEC v850" }, + { EM_M32R , "Renesas M32R (formerly Mitsubishi M32R)" }, + { EM_MN10300 , "Matsushita MN10300" }, + { EM_MN10200 , "Matsushita MN10200" }, + { EM_PJ , "picoJava" }, + { EM_OPENRISC , "OpenRISC 32-bit embedded processor" }, + { EM_ARC_A5 , "ARC Cores Tangent-A5" }, + { EM_XTENSA , "Tensilica Xtensa Architecture" }, + { EM_VIDEOCORE , "Alphamosaic VideoCore processor" }, + { EM_TMM_GPP , "Thompson Multimedia General Purpose Processor" }, + { EM_NS32K , "National Semiconductor 32000 series" }, + { EM_TPC , "Tenor Network TPC processor" }, + { EM_SNP1K , "Trebia SNP 1000 processor" }, + { EM_ST200 , "STMicroelectronics ST200 microcontroller" }, + { EM_IP2K , "Ubicom IP2022 micro controller" }, + { EM_MAX , "MAX Processor" }, + { EM_CR , "National Semiconductor CompactRISC" }, + { EM_F2MC16 , "Fujitsu F2MC16" }, + { EM_MSP430 , "TI msp430 micro controller" }, + { EM_BLACKFIN , "ADI Blackfin" }, + { EM_SE_C33 , "S1C33 Family of Seiko Epson processors" }, + { EM_SEP , "Sharp embedded microprocessor" }, + { EM_ARCA , "Arca RISC Microprocessor" }, + { EM_UNICORE , "Microprocessor series from PKU-Unity Ltd. and MPRC of Peking University" }, + { EM_EXCESS , "eXcess: 16/32/64-bit configurable embedded CPU" }, + { EM_DXP , "Icera Semiconductor Inc. Deep Execution Processor" }, + { EM_ALTERA_NIOS2 , "Altera Nios II soft-core processor" }, + { EM_CRX , "National Semiconductor CRX" }, + { EM_XGATE , "Motorola XGATE embedded processor" }, + { EM_C166 , "Infineon C16x/XC16x processor" }, + { EM_M16C , "Renesas M16C series microprocessors" }, + { EM_DSPIC30F , "Microchip Technology dsPIC30F Digital Signal Controller" }, + { EM_CE , "Freescale Communication Engine RISC core" }, + { EM_M32C , "Renesas M32C series microprocessors" }, + { EM_res121 , "Reserved" }, + { EM_res122 , "Reserved" }, + { EM_res123 , "Reserved" }, + { EM_res124 , "Reserved" }, + { EM_res125 , "Reserved" }, + { EM_res126 , "Reserved" }, + { EM_res127 , "Reserved" }, + { EM_res128 , "Reserved" }, + { EM_res129 , "Reserved" }, + { EM_res130 , "Reserved" }, + { EM_TSK3000 , "Altium TSK3000 core" }, + { EM_RS08 , "Freescale RS08 embedded processor" }, + { EM_res133 , "Reserved" }, + { EM_ECOG2 , "Cyan Technology eCOG2 microprocessor" }, + { EM_SCORE , "Sunplus Score" }, + { EM_SCORE7 , "Sunplus S+core7 RISC processor" }, + { EM_DSP24 , "New Japan Radio (NJR) 24-bit DSP Processor" }, + { EM_VIDEOCORE3 , "Broadcom VideoCore III processor" }, + { EM_LATTICEMICO32, "RISC processor for Lattice FPGA architecture" }, + { EM_SE_C17 , "Seiko Epson C17 family" }, + { EM_TI_C6000 , "Texas Instruments TMS320C6000 DSP family" }, + { EM_TI_C2000 , "Texas Instruments TMS320C2000 DSP family" }, + { EM_TI_C5500 , "Texas Instruments TMS320C55x DSP family" }, + { EM_res143 , "Reserved" }, + { EM_res144 , "Reserved" }, + { EM_res145 , "Reserved" }, + { EM_res146 , "Reserved" }, + { EM_res147 , "Reserved" }, + { EM_res148 , "Reserved" }, + { EM_res149 , "Reserved" }, + { EM_res150 , "Reserved" }, + { EM_res151 , "Reserved" }, + { EM_res152 , "Reserved" }, + { EM_res153 , "Reserved" }, + { EM_res154 , "Reserved" }, + { EM_res155 , "Reserved" }, + { EM_res156 , "Reserved" }, + { EM_res157 , "Reserved" }, + { EM_res158 , "Reserved" }, + { EM_res159 , "Reserved" }, + { EM_MMDSP_PLUS , "STMicroelectronics 64bit VLIW Data Signal Processor" }, + { EM_CYPRESS_M8C , "Cypress M8C microprocessor" }, + { EM_R32C , "Renesas R32C series microprocessors" }, + { EM_TRIMEDIA , "NXP Semiconductors TriMedia architecture family" }, + { EM_QDSP6 , "QUALCOMM DSP6 Processor" }, + { EM_8051 , "Intel 8051 and variants" }, + { EM_STXP7X , "STMicroelectronics STxP7x family" }, + { EM_NDS32 , "Andes Technology compact code size embedded RISC processor family" }, + { EM_ECOG1 , "Cyan Technology eCOG1X family" }, + { EM_ECOG1X , "Cyan Technology eCOG1X family" }, + { EM_MAXQ30 , "Dallas Semiconductor MAXQ30 Core Micro-controllers" }, + { EM_XIMO16 , "New Japan Radio (NJR) 16-bit DSP Processor" }, + { EM_MANIK , "M2000 Reconfigurable RISC Microprocessor" }, + { EM_CRAYNV2 , "Cray Inc. NV2 vector architecture" }, + { EM_RX , "Renesas RX family" }, + { EM_METAG , "Imagination Technologies META processor architecture" }, + { EM_MCST_ELBRUS , "MCST Elbrus general purpose hardware architecture" }, + { EM_ECOG16 , "Cyan Technology eCOG16 family" }, + { EM_CR16 , "National Semiconductor CompactRISC 16-bit processor" }, + { EM_ETPU , "Freescale Extended Time Processing Unit" }, + { EM_SLE9X , "Infineon Technologies SLE9X core" }, + { EM_L1OM , "Intel L1OM" }, + { EM_INTEL181 , "Reserved by Intel" }, + { EM_INTEL182 , "Reserved by Intel" }, + { EM_res183 , "Reserved by ARM" }, + { EM_res184 , "Reserved by ARM" }, + { EM_AVR32 , "Atmel Corporation 32-bit microprocessor family" }, + { EM_STM8 , "STMicroeletronics STM8 8-bit microcontroller" }, + { EM_TILE64 , "Tilera TILE64 multicore architecture family" }, + { EM_TILEPRO , "Tilera TILEPro multicore architecture family" }, + { EM_MICROBLAZE , "Xilinx MicroBlaze 32-bit RISC soft processor core" }, + { EM_CUDA , "NVIDIA CUDA architecture " }, +}; + + +static struct section_type_table_t { + const Elf64_Half key; + const char* str; +} section_type_table [] = +{ + { SHT_NULL , "NULL" }, + { SHT_PROGBITS , "PROGBITS" }, + { SHT_SYMTAB , "SYMTAB" }, + { SHT_STRTAB , "STRTAB" }, + { SHT_RELA , "RELA" }, + { SHT_HASH , "HASH" }, + { SHT_DYNAMIC , "DYNAMIC" }, + { SHT_NOTE , "NOTE" }, + { SHT_NOBITS , "NOBITS" }, + { SHT_REL , "REL" }, + { SHT_SHLIB , "SHLIB" }, + { SHT_DYNSYM , "DYNSYM" }, + { SHT_INIT_ARRAY , "INIT_ARRAY" }, + { SHT_FINI_ARRAY , "FINI_ARRAY" }, + { SHT_PREINIT_ARRAY, "PREINIT_ARRAY" }, + { SHT_GROUP , "GROUP" }, + { SHT_SYMTAB_SHNDX , "SYMTAB_SHNDX " }, +}; + + +static struct segment_type_table_t { + const Elf_Word key; + const char* str; +} segment_type_table [] = +{ + { PT_NULL , "NULL" }, + { PT_LOAD , "LOAD" }, + { PT_DYNAMIC, "DYNAMIC" }, + { PT_INTERP , "INTERP" }, + { PT_NOTE , "NOTE" }, + { PT_SHLIB , "SHLIB" }, + { PT_PHDR , "PHDR" }, + { PT_TLS , "TLS" }, +}; + + +static struct segment_flag_table_t { + const Elf_Word key; + const char* str; +} segment_flag_table [] = +{ + { 0, "" }, + { 1, "X" }, + { 2, "W" }, + { 3, "WX" }, + { 4, "R" }, + { 5, "RX" }, + { 6, "RW" }, + { 7, "RWX" }, +}; + + +static struct symbol_bind_t { + const Elf_Word key; + const char* str; +} symbol_bind_table [] = +{ + { STB_LOCAL , "LOCAL" }, + { STB_GLOBAL , "GLOBAL" }, + { STB_WEAK , "WEAK" }, + { STB_LOOS , "LOOS" }, + { STB_HIOS , "HIOS" }, + { STB_MULTIDEF, "MULTIDEF" }, + { STB_LOPROC , "LOPROC" }, + { STB_HIPROC , "HIPROC" }, +}; + + +static struct symbol_type_t { + const Elf_Word key; + const char* str; +} symbol_type_table [] = +{ + { STT_NOTYPE , "NOTYPE" }, + { STT_OBJECT , "OBJECT" }, + { STT_FUNC , "FUNC" }, + { STT_SECTION, "SECTION" }, + { STT_FILE , "FILE" }, + { STT_COMMON , "COMMON" }, + { STT_TLS , "TLS" }, + { STT_LOOS , "LOOS" }, + { STT_HIOS , "HIOS" }, + { STT_LOPROC , "LOPROC" }, + { STT_HIPROC , "HIPROC" }, +}; + + +static struct dynamic_tag_t { + const Elf_Word key; + const char* str; +} dynamic_tag_table [] = +{ + { DT_NULL , "NULL" }, + { DT_NEEDED , "NEEDED" }, + { DT_PLTRELSZ , "PLTRELSZ" }, + { DT_PLTGOT , "PLTGOT" }, + { DT_HASH , "HASH" }, + { DT_STRTAB , "STRTAB" }, + { DT_SYMTAB , "SYMTAB" }, + { DT_RELA , "RELA" }, + { DT_RELASZ , "RELASZ" }, + { DT_RELAENT , "RELAENT" }, + { DT_STRSZ , "STRSZ" }, + { DT_SYMENT , "SYMENT" }, + { DT_INIT , "INIT" }, + { DT_FINI , "FINI" }, + { DT_SONAME , "SONAME" }, + { DT_RPATH , "RPATH" }, + { DT_SYMBOLIC , "SYMBOLIC" }, + { DT_REL , "REL" }, + { DT_RELSZ , "RELSZ" }, + { DT_RELENT , "RELENT" }, + { DT_PLTREL , "PLTREL" }, + { DT_DEBUG , "DEBUG" }, + { DT_TEXTREL , "TEXTREL" }, + { DT_JMPREL , "JMPREL" }, + { DT_BIND_NOW , "BIND_NOW" }, + { DT_INIT_ARRAY , "INIT_ARRAY" }, + { DT_FINI_ARRAY , "FINI_ARRAY" }, + { DT_INIT_ARRAYSZ , "INIT_ARRAYSZ" }, + { DT_FINI_ARRAYSZ , "FINI_ARRAYSZ" }, + { DT_RUNPATH , "RUNPATH" }, + { DT_FLAGS , "FLAGS" }, + { DT_ENCODING , "ENCODING" }, + { DT_PREINIT_ARRAY , "PREINIT_ARRAY" }, + { DT_PREINIT_ARRAYSZ, "PREINIT_ARRAYSZ" }, + { DT_MAXPOSTAGS , "MAXPOSTAGS" }, +}; + +static const ELFIO::Elf_Xword MAX_DATA_ENTRIES = 64; + +//------------------------------------------------------------------------------ +class dump +{ +#define DUMP_DEC_FORMAT( width ) std::setw(width) << std::setfill( ' ' ) << \ + std::dec << std::right +#define DUMP_HEX_FORMAT( width ) std::setw(width) << std::setfill( '0' ) << \ + std::hex << std::right +#define DUMP_STR_FORMAT( width ) std::setw(width) << std::setfill( ' ' ) << \ + std::hex << std::left + + public: +//------------------------------------------------------------------------------ + static void + header( std::ostream& out, const elfio& reader ) + { + out << "ELF Header" << std::endl << std::endl + << " Class: " << str_class( reader.get_class() ) << std::endl + << " Encoding: " << str_endian( reader.get_encoding() ) << std::endl + << " ELFVersion: " << str_version( reader.get_elf_version() ) << std::endl + << " Type: " << str_type( reader.get_type() ) << std::endl + << " Machine: " << str_machine( reader.get_machine() ) << std::endl + << " Version: " << str_version( reader.get_version() ) << std::endl + << " Entry: " << "0x" << std::hex << reader.get_entry() << std::endl + << " Flags: " << "0x" << std::hex << reader.get_flags() << std::endl + << std::endl; + } + +//------------------------------------------------------------------------------ + static void + section_headers( std::ostream& out, const elfio& reader ) + { + Elf_Half n = reader.sections.size(); + + if ( n == 0 ) { + return; + } + + out << "Section Headers:" << std::endl; + if ( reader.get_class() == ELFCLASS32 ) { // Output for 32-bit + out << "[ Nr ] Type Addr Size ES Flg Lk Inf Al Name" << std::endl; + } + else { // Output for 64-bit + out << "[ Nr ] Type Addr Size ES Flg" << std::endl + << " Lk Inf Al Name" << std::endl; + } + + for ( Elf_Half i = 0; i < n; ++i ) { // For all sections + section* sec = reader.sections[i]; + section_header( out, i, sec, reader.get_class() ); + } + + out << "Key to Flags: W (write), A (alloc), X (execute)\n\n" + << std::endl; + } + +//------------------------------------------------------------------------------ + static void + section_header( std::ostream& out, Elf_Half no, const section* sec, + unsigned char elf_class ) + { + std::ios_base::fmtflags original_flags = out.flags(); + + if ( elf_class == ELFCLASS32 ) { // Output for 32-bit + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_STR_FORMAT( 17 ) << str_section_type( sec->get_type() ) << " " + << DUMP_HEX_FORMAT( 8 ) << sec->get_address() << " " + << DUMP_HEX_FORMAT( 8 ) << sec->get_size() << " " + << DUMP_HEX_FORMAT( 2 ) << sec->get_entry_size() << " " + << DUMP_STR_FORMAT( 3 ) << section_flags( sec->get_flags() ) << " " + << DUMP_HEX_FORMAT( 2 ) << sec->get_link() << " " + << DUMP_HEX_FORMAT( 3 ) << sec->get_info() << " " + << DUMP_HEX_FORMAT( 2 ) << sec->get_addr_align() << " " + << DUMP_STR_FORMAT( 17 ) << sec->get_name() << " " + << std::endl; + } + else { // Output for 64-bit + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_STR_FORMAT( 17 ) << str_section_type( sec->get_type() ) << " " + << DUMP_HEX_FORMAT( 16 ) << sec->get_address() << " " + << DUMP_HEX_FORMAT( 16 ) << sec->get_size() << " " + << DUMP_HEX_FORMAT( 4 ) << sec->get_entry_size() << " " + << DUMP_STR_FORMAT( 3 ) << section_flags( sec->get_flags() ) << " " + << std::endl + << " " + << DUMP_HEX_FORMAT( 4 ) << sec->get_link() << " " + << DUMP_HEX_FORMAT( 4 ) << sec->get_info() << " " + << DUMP_HEX_FORMAT( 4 ) << sec->get_addr_align() << " " + << DUMP_STR_FORMAT( 17 ) << sec->get_name() << " " + << std::endl; + } + + out.flags(original_flags); + + return; + } + +//------------------------------------------------------------------------------ + static void + segment_headers( std::ostream& out, const elfio& reader ) + { + Elf_Half n = reader.segments.size(); + if ( n == 0 ) { + return; + } + + out << "Segment headers:" << std::endl; + if ( reader.get_class() == ELFCLASS32 ) { // Output for 32-bit + out << "[ Nr ] Type VirtAddr PhysAddr FileSize Mem.Size Flags Align" + << std::endl; + } + else { // Output for 64-bit + out << "[ Nr ] Type VirtAddr PhysAddr Flags" << std::endl + << " FileSize Mem.Size Align" + << std::endl; + } + + for ( Elf_Half i = 0; i < n; ++i ) { + segment* seg = reader.segments[i]; + segment_header( out, i, seg, reader.get_class() ); + } + + out << std::endl; + } + +//------------------------------------------------------------------------------ + static void + segment_header( std::ostream& out, Elf_Half no, const segment* seg, + unsigned int elf_class ) + { + std::ios_base::fmtflags original_flags = out.flags(); + + if ( elf_class == ELFCLASS32 ) { // Output for 32-bit + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_STR_FORMAT( 14 ) << str_segment_type( seg->get_type() ) << " " + << DUMP_HEX_FORMAT( 8 ) << seg->get_virtual_address() << " " + << DUMP_HEX_FORMAT( 8 ) << seg->get_physical_address() << " " + << DUMP_HEX_FORMAT( 8 ) << seg->get_file_size() << " " + << DUMP_HEX_FORMAT( 8 ) << seg->get_memory_size() << " " + << DUMP_STR_FORMAT( 8 ) << str_segment_flag( seg->get_flags() ) << " " + << DUMP_HEX_FORMAT( 8 ) << seg->get_align() << " " + << std::endl; + } + else { // Output for 64-bit + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_STR_FORMAT( 14 ) << str_segment_type( seg->get_type() ) << " " + << DUMP_HEX_FORMAT( 16 ) << seg->get_virtual_address() << " " + << DUMP_HEX_FORMAT( 16 ) << seg->get_physical_address() << " " + << DUMP_STR_FORMAT( 16 ) << str_segment_flag( seg->get_flags() ) << " " + << std::endl + << " " + << DUMP_HEX_FORMAT( 16 ) << seg->get_file_size() << " " + << DUMP_HEX_FORMAT( 16 ) << seg->get_memory_size() << " " + << DUMP_HEX_FORMAT( 16 ) << seg->get_align() << " " + << std::endl; + } + + out.flags(original_flags); + } + +//------------------------------------------------------------------------------ + static void + symbol_tables( std::ostream& out, const elfio& reader ) + { + Elf_Half n = reader.sections.size(); + for ( Elf_Half i = 0; i < n; ++i ) { // For all sections + section* sec = reader.sections[i]; + if ( SHT_SYMTAB == sec->get_type() || SHT_DYNSYM == sec->get_type() ) { + symbol_section_accessor symbols( reader, sec ); + + Elf_Xword sym_no = symbols.get_symbols_num(); + if ( sym_no > 0 ) { + out << "Symbol table (" << sec->get_name() << ")" << std::endl; + if ( reader.get_class() == ELFCLASS32 ) { // Output for 32-bit + out << "[ Nr ] Value Size Type Bind Sect Name" + << std::endl; + } + else { // Output for 64-bit + out << "[ Nr ] Value Size Type Bind Sect" << std::endl + << " Name" + << std::endl; + } + for ( Elf_Half i = 0; i < sym_no; ++i ) { + std::string name; + Elf64_Addr value = 0; + Elf_Xword size = 0; + unsigned char bind = 0; + unsigned char type = 0; + Elf_Half section = 0; + unsigned char other = 0; + symbols.get_symbol( i, name, value, size, bind, type, section, other ); + symbol_table( out, i, name, value, size, bind, type, section, reader.get_class() ); + } + + out << std::endl; + } + } + } + } + +//------------------------------------------------------------------------------ + static void + symbol_table( std::ostream& out, + Elf_Half no, + std::string& name, + Elf64_Addr value, + Elf_Xword size, + unsigned char bind, + unsigned char type, + Elf_Half section, + unsigned int elf_class ) + { + std::ios_base::fmtflags original_flags = out.flags(); + + if ( elf_class == ELFCLASS32 ) { // Output for 32-bit + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_HEX_FORMAT( 8 ) << value << " " + << DUMP_HEX_FORMAT( 8 ) << size << " " + << DUMP_STR_FORMAT( 7 ) << str_symbol_type( type ) << " " + << DUMP_STR_FORMAT( 8 ) << str_symbol_bind( bind ) << " " + << DUMP_DEC_FORMAT( 5 ) << section << " " + << DUMP_STR_FORMAT( 1 ) << name << " " + << std::endl; + } + else { // Output for 64-bit + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_HEX_FORMAT( 16 ) << value << " " + << DUMP_HEX_FORMAT( 16 ) << size << " " + << DUMP_STR_FORMAT( 7 ) << str_symbol_type( type ) << " " + << DUMP_STR_FORMAT( 8 ) << str_symbol_bind( bind ) << " " + << DUMP_DEC_FORMAT( 5 ) << section << " " + << std::endl + << " " + << DUMP_STR_FORMAT( 1 ) << name << " " + << std::endl; + } + + out.flags(original_flags); + } + +//------------------------------------------------------------------------------ + static void + notes( std::ostream& out, const elfio& reader ) + { + Elf_Half no = reader.sections.size(); + for ( Elf_Half i = 0; i < no; ++i ) { // For all sections + section* sec = reader.sections[i]; + if ( SHT_NOTE == sec->get_type() ) { // Look at notes + note_section_accessor notes( reader, sec ); + int no_notes = notes.get_notes_num(); + if ( no > 0 ) { + out << "Note section (" << sec->get_name() << ")" << std::endl + << " No Type Name" + << std::endl; + for ( int j = 0; j < no_notes; ++j ) { // For all notes + Elf_Word type; + std::string name; + void* desc; + Elf_Word descsz; + + if ( notes.get_note(j, type, name, desc, descsz) ) { + // 'name' usually contains \0 at the end. Try to fix it + name = name.c_str(); + note( out, j, type, name ); + } + } + + out << std::endl; + } + } + } + } + +//------------------------------------------------------------------------------ + static void + note( std::ostream& out, + int no, + Elf_Word type, + const std::string& name ) + { + out << " [" + << DUMP_DEC_FORMAT( 2 ) << no + << "] " + << DUMP_HEX_FORMAT( 8 ) << type << " " + << DUMP_STR_FORMAT( 1 ) << name + << std::endl; + } + +//------------------------------------------------------------------------------ + static void + dynamic_tags( std::ostream& out, const elfio& reader ) + { + Elf_Half n = reader.sections.size(); + for ( Elf_Half i = 0; i < n; ++i ) { // For all sections + section* sec = reader.sections[i]; + if ( SHT_DYNAMIC == sec->get_type() ) { + dynamic_section_accessor dynamic( reader, sec ); + + Elf_Xword dyn_no = dynamic.get_entries_num(); + if ( dyn_no > 0 ) { + out << "Dynamic section (" << sec->get_name() << ")" << std::endl; + out << "[ Nr ] Tag Name/Value" << std::endl; + for ( int i = 0; i < dyn_no; ++i ) { + Elf_Xword tag = 0; + Elf_Xword value = 0; + std::string str; + dynamic.get_entry( i, tag, value, str ); + dynamic_tag( out, i, tag, value, str, reader.get_class() ); + if ( DT_NULL == tag ) { + break; + } + } + + out << std::endl; + } + } + } + } + +//------------------------------------------------------------------------------ + static void + dynamic_tag( std::ostream& out, + int no, + Elf_Xword tag, + Elf_Xword value, + std::string str, + unsigned int /*elf_class*/ ) + { + out << "[" + << DUMP_DEC_FORMAT( 5 ) << no + << "] " + << DUMP_STR_FORMAT( 16 ) << str_dynamic_tag( tag ) << " "; + if ( str.empty() ) { + out << DUMP_HEX_FORMAT( 16 ) << value << " "; + } + else { + out << DUMP_STR_FORMAT( 32 ) << str << " "; + } + out << std::endl; + } + +//------------------------------------------------------------------------------ + static void + section_data( std::ostream& out, const section* sec ) + { + std::ios_base::fmtflags original_flags = out.flags(); + + out << sec->get_name() << std::endl; + const char* pdata = sec->get_data(); + if ( pdata ){ + ELFIO::Elf_Xword i; + for ( i = 0; i < std::min( sec->get_size(), MAX_DATA_ENTRIES ); ++i ) { + if ( i % 16 == 0 ) { + out << "[" << DUMP_HEX_FORMAT( 8 ) << i << "]"; + } + + out << " " << DUMP_HEX_FORMAT( 2 ) << ( pdata[i] & 0x000000FF ); + + if ( i % 16 == 15 ) { + out << std::endl; + } + } + if ( i % 16 != 0 ) { + out << std::endl; + } + + out.flags(original_flags); + } + + return; + } + +//------------------------------------------------------------------------------ + static void + section_datas( std::ostream& out, const elfio& reader ) + { + Elf_Half n = reader.sections.size(); + + if ( n == 0 ) { + return; + } + + out << "Section Data:" << std::endl; + + for ( Elf_Half i = 1; i < n; ++i ) { // For all sections + section* sec = reader.sections[i]; + if ( sec->get_type() == SHT_NOBITS ) { + continue; + } + section_data( out, sec ); + } + + out << std::endl; + } + +//------------------------------------------------------------------------------ + static void + segment_data( std::ostream& out, Elf_Half no, const segment* seg ) + { + std::ios_base::fmtflags original_flags = out.flags(); + + out << "Segment # " << no << std::endl; + const char* pdata = seg->get_data(); + if ( pdata ) { + ELFIO::Elf_Xword i; + for ( i = 0; i < std::min( seg->get_file_size(), MAX_DATA_ENTRIES ); ++i ) { + if ( i % 16 == 0 ) { + out << "[" << DUMP_HEX_FORMAT( 8 ) << i << "]"; + } + + out << " " << DUMP_HEX_FORMAT( 2 ) << ( pdata[i] & 0x000000FF ); + + if ( i % 16 == 15 ) { + out << std::endl; + } + } + if ( i % 16 != 0 ) { + out << std::endl; + } + + out.flags(original_flags); + } + + return; + } + +//------------------------------------------------------------------------------ + static void + segment_datas( std::ostream& out, const elfio& reader ) + { + Elf_Half n = reader.segments.size(); + + if ( n == 0 ) { + return; + } + + out << "Segment Data:" << std::endl; + + for ( Elf_Half i = 0; i < n; ++i ) { // For all sections + segment* seg = reader.segments[i]; + segment_data( out, i, seg ); + } + + out << std::endl; + } + + private: +//------------------------------------------------------------------------------ + template< typename T, typename K > + std::string + static + find_value_in_table( const T& table, const K& key ) + { + std::string res = "?"; + for ( unsigned int i = 0; i < sizeof( table )/sizeof( table[0] ); ++i ) { + if ( table[i].key == key ) { + res = table[i].str; + break; + } + } + + return res; + } + + +//------------------------------------------------------------------------------ + template< typename T, typename K > + static + std::string + format_assoc( const T& table, const K& key ) + { + std::string str = find_value_in_table( table, key ); + if ( str == "?" ) { + std::ostringstream oss; + oss << str << " (0x" << std::hex << key << ")"; + str = oss.str(); + } + + return str; + } + + +//------------------------------------------------------------------------------ + template< typename T > + static + std::string + format_assoc( const T& table, const char key ) + { + return format_assoc( table, (const int)key ); + } + + +//------------------------------------------------------------------------------ + static + std::string + section_flags( Elf_Xword flags ) + { + std::string ret = ""; + if ( flags & SHF_WRITE ) { + ret += "W"; + } + if ( flags & SHF_ALLOC ) { + ret += "A"; + } + if ( flags & SHF_EXECINSTR ) { + ret += "X"; + } + + return ret; + } + + +//------------------------------------------------------------------------------ +#define STR_FUNC_TABLE( name ) \ + template< typename T > \ + static \ + std::string \ + str_##name( const T key ) \ + { \ + return format_assoc( name##_table, key ); \ + } + + STR_FUNC_TABLE( class ) + STR_FUNC_TABLE( endian ) + STR_FUNC_TABLE( version ) + STR_FUNC_TABLE( type ) + STR_FUNC_TABLE( machine ) + STR_FUNC_TABLE( section_type ) + STR_FUNC_TABLE( segment_type ) + STR_FUNC_TABLE( segment_flag ) + STR_FUNC_TABLE( symbol_bind ) + STR_FUNC_TABLE( symbol_type ) + STR_FUNC_TABLE( dynamic_tag ) + +#undef STR_FUNC_TABLE +#undef DUMP_DEC_FORMAT +#undef DUMP_HEX_FORMAT +#undef DUMP_STR_FORMAT +}; // class dump + + +}; // namespace ELFIO + +#endif // ELFIO_DUMP_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dynamic.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dynamic.hpp new file mode 100644 index 0000000000..6f2d041e0f --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dynamic.hpp @@ -0,0 +1,253 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_DYNAMIC_HPP +#define ELFIO_DYNAMIC_HPP + +namespace ELFIO { + +//------------------------------------------------------------------------------ +class dynamic_section_accessor +{ + public: +//------------------------------------------------------------------------------ + dynamic_section_accessor( const elfio& elf_file_, section* section_ ) : + elf_file( elf_file_ ), + dynamic_section( section_ ) + { + } + +//------------------------------------------------------------------------------ + Elf_Xword + get_entries_num() const + { + Elf_Xword nRet = 0; + + if ( 0 != dynamic_section->get_entry_size() ) { + nRet = dynamic_section->get_size() / dynamic_section->get_entry_size(); + } + + return nRet; + } + +//------------------------------------------------------------------------------ + bool + get_entry( Elf_Xword index, + Elf_Xword& tag, + Elf_Xword& value, + std::string& str ) const + { + if ( index >= get_entries_num() ) { // Is index valid + return false; + } + + if ( elf_file.get_class() == ELFCLASS32 ) { + generic_get_entry_dyn< Elf32_Dyn >( index, tag, value ); + } + else { + generic_get_entry_dyn< Elf64_Dyn >( index, tag, value ); + } + + // If the tag may have a string table reference, prepare the string + if ( tag == DT_NEEDED || + tag == DT_SONAME || + tag == DT_RPATH || + tag == DT_RUNPATH ) { + string_section_accessor strsec = + elf_file.sections[ get_string_table_index() ]; + const char* result = strsec.get_string( value ); + if ( 0 == result ) { + str.clear(); + return false; + } + str = result; + } + else { + str.clear(); + } + + return true; + } + +//------------------------------------------------------------------------------ + void + add_entry( Elf_Xword& tag, + Elf_Xword& value ) + { + if ( elf_file.get_class() == ELFCLASS32 ) { + generic_add_entry< Elf32_Dyn >( tag, value ); + } + else { + generic_add_entry< Elf64_Dyn >( tag, value ); + } + } + +//------------------------------------------------------------------------------ + void + add_entry( Elf_Xword& tag, + std::string& str ) + { + string_section_accessor strsec = + elf_file.sections[ get_string_table_index() ]; + Elf_Xword value = strsec.add_string( str ); + add_entry( tag, value ); + } + +//------------------------------------------------------------------------------ + private: +//------------------------------------------------------------------------------ + Elf_Half + get_string_table_index() const + { + return (Elf_Half)dynamic_section->get_link(); + } + +//------------------------------------------------------------------------------ + template< class T > + void + generic_get_entry_dyn( Elf_Xword index, + Elf_Xword& tag, + Elf_Xword& value ) const + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + // Check unusual case when dynamic section has no data + if( dynamic_section->get_data() == 0 || + ( index + 1 ) * dynamic_section->get_entry_size() > dynamic_section->get_size() ) { + tag = DT_NULL; + value = 0; + return; + } + + const T* pEntry = reinterpret_cast( + dynamic_section->get_data() + + index * dynamic_section->get_entry_size() ); + tag = convertor( pEntry->d_tag ); + switch ( tag ) { + case DT_NULL: + case DT_SYMBOLIC: + case DT_TEXTREL: + case DT_BIND_NOW: + value = 0; + break; + case DT_NEEDED: + case DT_PLTRELSZ: + case DT_RELASZ: + case DT_RELAENT: + case DT_STRSZ: + case DT_SYMENT: + case DT_SONAME: + case DT_RPATH: + case DT_RELSZ: + case DT_RELENT: + case DT_PLTREL: + case DT_INIT_ARRAYSZ: + case DT_FINI_ARRAYSZ: + case DT_RUNPATH: + case DT_FLAGS: + case DT_PREINIT_ARRAYSZ: + value = convertor( pEntry->d_un.d_val ); + break; + case DT_PLTGOT: + case DT_HASH: + case DT_STRTAB: + case DT_SYMTAB: + case DT_RELA: + case DT_INIT: + case DT_FINI: + case DT_REL: + case DT_DEBUG: + case DT_JMPREL: + case DT_INIT_ARRAY: + case DT_FINI_ARRAY: + case DT_PREINIT_ARRAY: + default: + value = convertor( pEntry->d_un.d_ptr ); + break; + } + } + +//------------------------------------------------------------------------------ + template< class T > + void + generic_add_entry( Elf_Xword tag, Elf_Xword value ) + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + T entry; + + switch ( tag ) { + case DT_NULL: + case DT_SYMBOLIC: + case DT_TEXTREL: + case DT_BIND_NOW: + value = 0; + case DT_NEEDED: + case DT_PLTRELSZ: + case DT_RELASZ: + case DT_RELAENT: + case DT_STRSZ: + case DT_SYMENT: + case DT_SONAME: + case DT_RPATH: + case DT_RELSZ: + case DT_RELENT: + case DT_PLTREL: + case DT_INIT_ARRAYSZ: + case DT_FINI_ARRAYSZ: + case DT_RUNPATH: + case DT_FLAGS: + case DT_PREINIT_ARRAYSZ: + entry.d_un.d_val = convertor( value ); + break; + case DT_PLTGOT: + case DT_HASH: + case DT_STRTAB: + case DT_SYMTAB: + case DT_RELA: + case DT_INIT: + case DT_FINI: + case DT_REL: + case DT_DEBUG: + case DT_JMPREL: + case DT_INIT_ARRAY: + case DT_FINI_ARRAY: + case DT_PREINIT_ARRAY: + default: + entry.d_un.d_ptr = convertor( value ); + break; + } + + entry.d_tag = convertor( tag ); + + dynamic_section->append_data( reinterpret_cast( &entry ), sizeof( entry ) ); + } + +//------------------------------------------------------------------------------ + private: + const elfio& elf_file; + section* dynamic_section; +}; + +} // namespace ELFIO + +#endif // ELFIO_DYNAMIC_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_header.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_header.hpp new file mode 100644 index 0000000000..d689a8899f --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_header.hpp @@ -0,0 +1,146 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELF_HEADER_HPP +#define ELF_HEADER_HPP + +#include + +namespace ELFIO { + +class elf_header +{ + public: + virtual ~elf_header() {}; + virtual bool load( std::istream& stream ) = 0; + virtual bool save( std::ostream& stream ) const = 0; + + // ELF header functions + ELFIO_GET_ACCESS_DECL( unsigned char, class ); + ELFIO_GET_ACCESS_DECL( unsigned char, elf_version ); + ELFIO_GET_ACCESS_DECL( unsigned char, encoding ); + ELFIO_GET_ACCESS_DECL( Elf_Word, version ); + ELFIO_GET_ACCESS_DECL( Elf_Half, header_size ); + ELFIO_GET_ACCESS_DECL( Elf_Half, section_entry_size ); + ELFIO_GET_ACCESS_DECL( Elf_Half, segment_entry_size ); + + ELFIO_GET_SET_ACCESS_DECL( unsigned char, os_abi ); + ELFIO_GET_SET_ACCESS_DECL( unsigned char, abi_version ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Half, type ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Half, machine ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, flags ); + ELFIO_GET_SET_ACCESS_DECL( Elf64_Addr, entry ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Half, sections_num ); + ELFIO_GET_SET_ACCESS_DECL( Elf64_Off, sections_offset ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Half, segments_num ); + ELFIO_GET_SET_ACCESS_DECL( Elf64_Off, segments_offset ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Half, section_name_str_index ); +}; + + +template< class T > struct elf_header_impl_types; +template<> struct elf_header_impl_types { + typedef Elf32_Phdr Phdr_type; + typedef Elf32_Shdr Shdr_type; + static const unsigned char file_class = ELFCLASS32; +}; +template<> struct elf_header_impl_types { + typedef Elf64_Phdr Phdr_type; + typedef Elf64_Shdr Shdr_type; + static const unsigned char file_class = ELFCLASS64; +}; + +template< class T > class elf_header_impl : public elf_header +{ + public: + elf_header_impl( endianess_convertor* convertor_, + unsigned char encoding ) + { + convertor = convertor_; + + std::fill_n( reinterpret_cast( &header ), sizeof( header ), '\0' ); + + header.e_ident[EI_MAG0] = ELFMAG0; + header.e_ident[EI_MAG1] = ELFMAG1; + header.e_ident[EI_MAG2] = ELFMAG2; + header.e_ident[EI_MAG3] = ELFMAG3; + header.e_ident[EI_CLASS] = elf_header_impl_types::file_class; + header.e_ident[EI_DATA] = encoding; + header.e_ident[EI_VERSION] = EV_CURRENT; + header.e_version = EV_CURRENT; + header.e_version = (*convertor)( header.e_version ); + header.e_ehsize = ( sizeof( header ) ); + header.e_ehsize = (*convertor)( header.e_ehsize ); + header.e_shstrndx = (*convertor)( (Elf_Half)1 ); + header.e_phentsize = sizeof( typename elf_header_impl_types::Phdr_type ); + header.e_shentsize = sizeof( typename elf_header_impl_types::Shdr_type ); + header.e_phentsize = (*convertor)( header.e_phentsize ); + header.e_shentsize = (*convertor)( header.e_shentsize ); + } + + bool + load( std::istream& stream ) + { + stream.seekg( 0 ); + stream.read( reinterpret_cast( &header ), sizeof( header ) ); + + return (stream.gcount() == sizeof( header ) ); + } + + bool + save( std::ostream& stream ) const + { + stream.seekp( 0 ); + stream.write( reinterpret_cast( &header ), sizeof( header ) ); + + return stream.good(); + } + + // ELF header functions + ELFIO_GET_ACCESS( unsigned char, class, header.e_ident[EI_CLASS] ); + ELFIO_GET_ACCESS( unsigned char, elf_version, header.e_ident[EI_VERSION] ); + ELFIO_GET_ACCESS( unsigned char, encoding, header.e_ident[EI_DATA] ); + ELFIO_GET_ACCESS( Elf_Word, version, header.e_version ); + ELFIO_GET_ACCESS( Elf_Half, header_size, header.e_ehsize ); + ELFIO_GET_ACCESS( Elf_Half, section_entry_size, header.e_shentsize ); + ELFIO_GET_ACCESS( Elf_Half, segment_entry_size, header.e_phentsize ); + + ELFIO_GET_SET_ACCESS( unsigned char, os_abi, header.e_ident[EI_OSABI] ); + ELFIO_GET_SET_ACCESS( unsigned char, abi_version, header.e_ident[EI_ABIVERSION] ); + ELFIO_GET_SET_ACCESS( Elf_Half, type, header.e_type ); + ELFIO_GET_SET_ACCESS( Elf_Half, machine, header.e_machine ); + ELFIO_GET_SET_ACCESS( Elf_Word, flags, header.e_flags ); + ELFIO_GET_SET_ACCESS( Elf_Half, section_name_str_index, header.e_shstrndx ); + ELFIO_GET_SET_ACCESS( Elf64_Addr, entry, header.e_entry ); + ELFIO_GET_SET_ACCESS( Elf_Half, sections_num, header.e_shnum ); + ELFIO_GET_SET_ACCESS( Elf64_Off, sections_offset, header.e_shoff ); + ELFIO_GET_SET_ACCESS( Elf_Half, segments_num, header.e_phnum ); + ELFIO_GET_SET_ACCESS( Elf64_Off, segments_offset, header.e_phoff ); + + private: + T header; + endianess_convertor* convertor; +}; + +} // namespace ELFIO + +#endif // ELF_HEADER_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_note.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_note.hpp new file mode 100644 index 0000000000..35c6fe344c --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_note.hpp @@ -0,0 +1,166 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_NOTE_HPP +#define ELFIO_NOTE_HPP + +namespace ELFIO { + +//------------------------------------------------------------------------------ +// There are discrepancies in documentations. SCO documentation +// (http://www.sco.com/developers/gabi/latest/ch5.pheader.html#note_section) +// requires 8 byte entries alignment for 64-bit ELF file, +// but Oracle's definition uses the same structure +// for 32-bit and 64-bit formats. +// (https://docs.oracle.com/cd/E23824_01/html/819-0690/chapter6-18048.html) +// +// It looks like EM_X86_64 Linux implementation is similar to Oracle's +// definition. Therefore, the same alignment works for both formats +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ +class note_section_accessor +{ + public: +//------------------------------------------------------------------------------ + note_section_accessor( const elfio& elf_file_, section* section_ ) : + elf_file( elf_file_ ), note_section( section_ ) + { + process_section(); + } + +//------------------------------------------------------------------------------ + Elf_Word + get_notes_num() const + { + return (Elf_Word)note_start_positions.size(); + } + +//------------------------------------------------------------------------------ + bool + get_note( Elf_Word index, + Elf_Word& type, + std::string& name, + void*& desc, + Elf_Word& descSize ) const + { + if ( index >= note_section->get_size() ) { + return false; + } + + const char* pData = note_section->get_data() + note_start_positions[index]; + int align = sizeof( Elf_Word ); + + const endianess_convertor& convertor = elf_file.get_convertor(); + type = convertor( *(Elf_Word*)( pData + 2*align ) ); + Elf_Word namesz = convertor( *(Elf_Word*)( pData ) ); + descSize = convertor( *(Elf_Word*)( pData + sizeof( namesz ) ) ); + Elf_Word max_name_size = note_section->get_size() - note_start_positions[index]; + if ( namesz > max_name_size || + namesz + descSize > max_name_size ) { + return false; + } + name.assign( pData + 3*align, namesz - 1); + if ( 0 == descSize ) { + desc = 0; + } + else { + desc = const_cast ( pData + 3*align + + ( ( namesz + align - 1 )/align )*align ); + } + + return true; + } + +//------------------------------------------------------------------------------ + void add_note( Elf_Word type, + const std::string& name, + const void* desc, + Elf_Word descSize ) + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + int align = sizeof( Elf_Word ); + Elf_Word nameLen = (Elf_Word)name.size() + 1; + Elf_Word nameLenConv = convertor( nameLen ); + std::string buffer( reinterpret_cast( &nameLenConv ), align ); + Elf_Word descSizeConv = convertor( descSize ); + buffer.append( reinterpret_cast( &descSizeConv ), align ); + type = convertor( type ); + buffer.append( reinterpret_cast( &type ), align ); + buffer.append( name ); + buffer.append( 1, '\x00' ); + const char pad[] = { '\0', '\0', '\0', '\0' }; + if ( nameLen % align != 0 ) { + buffer.append( pad, align - nameLen % align ); + } + if ( desc != 0 && descSize != 0 ) { + buffer.append( reinterpret_cast( desc ), descSize ); + if ( descSize % align != 0 ) { + buffer.append( pad, align - descSize % align ); + } + } + + note_start_positions.push_back( note_section->get_size() ); + note_section->append_data( buffer ); + } + + private: +//------------------------------------------------------------------------------ + void process_section() + { + const endianess_convertor& convertor = elf_file.get_convertor(); + const char* data = note_section->get_data(); + Elf_Xword size = note_section->get_size(); + Elf_Xword current = 0; + + note_start_positions.clear(); + + // Is it empty? + if ( 0 == data || 0 == size ) { + return; + } + + int align = sizeof( Elf_Word ); + while ( current + 3*align <= size ) { + note_start_positions.push_back( current ); + Elf_Word namesz = convertor( + *(Elf_Word*)( data + current ) ); + Elf_Word descsz = convertor( + *(Elf_Word*)( data + current + sizeof( namesz ) ) ); + + current += 3*sizeof( Elf_Word ) + + ( ( namesz + align - 1 ) / align ) * align + + ( ( descsz + align - 1 ) / align ) * align; + } + } + +//------------------------------------------------------------------------------ + private: + const elfio& elf_file; + section* note_section; + std::vector note_start_positions; +}; + +} // namespace ELFIO + +#endif // ELFIO_NOTE_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_relocation.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_relocation.hpp new file mode 100644 index 0000000000..d13d8b23c7 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_relocation.hpp @@ -0,0 +1,369 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_RELOCATION_HPP +#define ELFIO_RELOCATION_HPP + +namespace ELFIO { + +template struct get_sym_and_type; +template<> struct get_sym_and_type< Elf32_Rel > +{ + static int get_r_sym( Elf_Xword info ) + { + return ELF32_R_SYM( (Elf_Word)info ); + } + static int get_r_type( Elf_Xword info ) + { + return ELF32_R_TYPE( (Elf_Word)info ); + } +}; +template<> struct get_sym_and_type< Elf32_Rela > +{ + static int get_r_sym( Elf_Xword info ) + { + return ELF32_R_SYM( (Elf_Word)info ); + } + static int get_r_type( Elf_Xword info ) + { + return ELF32_R_TYPE( (Elf_Word)info ); + } +}; +template<> struct get_sym_and_type< Elf64_Rel > +{ + static int get_r_sym( Elf_Xword info ) + { + return ELF64_R_SYM( info ); + } + static int get_r_type( Elf_Xword info ) + { + return ELF64_R_TYPE( info ); + } +}; +template<> struct get_sym_and_type< Elf64_Rela > +{ + static int get_r_sym( Elf_Xword info ) + { + return ELF64_R_SYM( info ); + } + static int get_r_type( Elf_Xword info ) + { + return ELF64_R_TYPE( info ); + } +}; + + +//------------------------------------------------------------------------------ +class relocation_section_accessor +{ + public: +//------------------------------------------------------------------------------ + relocation_section_accessor( const elfio& elf_file_, section* section_ ) : + elf_file( elf_file_ ), + relocation_section( section_ ) + { + } + +//------------------------------------------------------------------------------ + Elf_Xword + get_entries_num() const + { + Elf_Xword nRet = 0; + + if ( 0 != relocation_section->get_entry_size() ) { + nRet = relocation_section->get_size() / relocation_section->get_entry_size(); + } + + return nRet; + } + +//------------------------------------------------------------------------------ + bool + get_entry( Elf_Xword index, + Elf64_Addr& offset, + Elf_Word& symbol, + Elf_Word& type, + Elf_Sxword& addend ) const + { + if ( index >= get_entries_num() ) { // Is index valid + return false; + } + + if ( elf_file.get_class() == ELFCLASS32 ) { + if ( SHT_REL == relocation_section->get_type() ) { + generic_get_entry_rel< Elf32_Rel >( index, offset, symbol, + type, addend ); + } + else if ( SHT_RELA == relocation_section->get_type() ) { + generic_get_entry_rela< Elf32_Rela >( index, offset, symbol, + type, addend ); + } + } + else { + if ( SHT_REL == relocation_section->get_type() ) { + generic_get_entry_rel< Elf64_Rel >( index, offset, symbol, + type, addend ); + } + else if ( SHT_RELA == relocation_section->get_type() ) { + generic_get_entry_rela< Elf64_Rela >( index, offset, symbol, + type, addend ); + } + } + + return true; + } + +//------------------------------------------------------------------------------ + bool + get_entry( Elf_Xword index, + Elf64_Addr& offset, + Elf64_Addr& symbolValue, + std::string& symbolName, + Elf_Word& type, + Elf_Sxword& addend, + Elf_Sxword& calcValue ) const + { + // Do regular job + Elf_Word symbol; + bool ret = get_entry( index, offset, symbol, type, addend ); + + // Find the symbol + Elf_Xword size; + unsigned char bind; + unsigned char symbolType; + Elf_Half section; + unsigned char other; + + symbol_section_accessor symbols( elf_file, elf_file.sections[get_symbol_table_index()] ); + ret = ret && symbols.get_symbol( symbol, symbolName, symbolValue, + size, bind, symbolType, section, other ); + + if ( ret ) { // Was it successful? + switch ( type ) { + case R_386_NONE: // none + calcValue = 0; + break; + case R_386_32: // S + A + calcValue = symbolValue + addend; + break; + case R_386_PC32: // S + A - P + calcValue = symbolValue + addend - offset; + break; + case R_386_GOT32: // G + A - P + calcValue = 0; + break; + case R_386_PLT32: // L + A - P + calcValue = 0; + break; + case R_386_COPY: // none + calcValue = 0; + break; + case R_386_GLOB_DAT: // S + case R_386_JMP_SLOT: // S + calcValue = symbolValue; + break; + case R_386_RELATIVE: // B + A + calcValue = addend; + break; + case R_386_GOTOFF: // S + A - GOT + calcValue = 0; + break; + case R_386_GOTPC: // GOT + A - P + calcValue = 0; + break; + default: // Not recognized symbol! + calcValue = 0; + break; + } + } + + return ret; + } + +//------------------------------------------------------------------------------ + void + add_entry( Elf64_Addr offset, Elf_Xword info ) + { + if ( elf_file.get_class() == ELFCLASS32 ) { + generic_add_entry< Elf32_Rel >( offset, info ); + } + else { + generic_add_entry< Elf64_Rel >( offset, info ); + } + } + +//------------------------------------------------------------------------------ + void + add_entry( Elf64_Addr offset, Elf_Word symbol, unsigned char type ) + { + Elf_Xword info; + if ( elf_file.get_class() == ELFCLASS32 ) { + info = ELF32_R_INFO( (Elf_Xword)symbol, type ); + } + else { + info = ELF64_R_INFO((Elf_Xword)symbol, type ); + } + + add_entry( offset, info ); + } + +//------------------------------------------------------------------------------ + void + add_entry( Elf64_Addr offset, Elf_Xword info, Elf_Sxword addend ) + { + if ( elf_file.get_class() == ELFCLASS32 ) { + generic_add_entry< Elf32_Rela >( offset, info, addend ); + } + else { + generic_add_entry< Elf64_Rela >( offset, info, addend ); + } + } + +//------------------------------------------------------------------------------ + void + add_entry( Elf64_Addr offset, Elf_Word symbol, unsigned char type, + Elf_Sxword addend ) + { + Elf_Xword info; + if ( elf_file.get_class() == ELFCLASS32 ) { + info = ELF32_R_INFO( (Elf_Xword)symbol, type ); + } + else { + info = ELF64_R_INFO( (Elf_Xword)symbol, type ); + } + + add_entry( offset, info, addend ); + } + +//------------------------------------------------------------------------------ + void + add_entry( string_section_accessor str_writer, + const char* str, + symbol_section_accessor sym_writer, + Elf64_Addr value, + Elf_Word size, + unsigned char sym_info, + unsigned char other, + Elf_Half shndx, + Elf64_Addr offset, + unsigned char type ) + { + Elf_Word str_index = str_writer.add_string( str ); + Elf_Word sym_index = sym_writer.add_symbol( str_index, value, size, + sym_info, other, shndx ); + add_entry( offset, sym_index, type ); + } + +//------------------------------------------------------------------------------ + private: +//------------------------------------------------------------------------------ + Elf_Half + get_symbol_table_index() const + { + return (Elf_Half)relocation_section->get_link(); + } + +//------------------------------------------------------------------------------ + template< class T > + void + generic_get_entry_rel( Elf_Xword index, + Elf64_Addr& offset, + Elf_Word& symbol, + Elf_Word& type, + Elf_Sxword& addend ) const + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + const T* pEntry = reinterpret_cast( + relocation_section->get_data() + + index * relocation_section->get_entry_size() ); + offset = convertor( pEntry->r_offset ); + Elf_Xword tmp = convertor( pEntry->r_info ); + symbol = get_sym_and_type::get_r_sym( tmp ); + type = get_sym_and_type::get_r_type( tmp ); + addend = 0; + } + +//------------------------------------------------------------------------------ + template< class T > + void + generic_get_entry_rela( Elf_Xword index, + Elf64_Addr& offset, + Elf_Word& symbol, + Elf_Word& type, + Elf_Sxword& addend ) const + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + const T* pEntry = reinterpret_cast( + relocation_section->get_data() + + index * relocation_section->get_entry_size() ); + offset = convertor( pEntry->r_offset ); + Elf_Xword tmp = convertor( pEntry->r_info ); + symbol = get_sym_and_type::get_r_sym( tmp ); + type = get_sym_and_type::get_r_type( tmp ); + addend = convertor( pEntry->r_addend ); + } + +//------------------------------------------------------------------------------ + template< class T > + void + generic_add_entry( Elf64_Addr offset, Elf_Xword info ) + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + T entry; + entry.r_offset = offset; + entry.r_info = info; + entry.r_offset = convertor( entry.r_offset ); + entry.r_info = convertor( entry.r_info ); + + relocation_section->append_data( reinterpret_cast( &entry ), sizeof( entry ) ); + } + +//------------------------------------------------------------------------------ + template< class T > + void + generic_add_entry( Elf64_Addr offset, Elf_Xword info, Elf_Sxword addend ) + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + T entry; + entry.r_offset = offset; + entry.r_info = info; + entry.r_addend = addend; + entry.r_offset = convertor( entry.r_offset ); + entry.r_info = convertor( entry.r_info ); + entry.r_addend = convertor( entry.r_addend ); + + relocation_section->append_data( reinterpret_cast( &entry ), sizeof( entry ) ); + } + +//------------------------------------------------------------------------------ + private: + const elfio& elf_file; + section* relocation_section; +}; + +} // namespace ELFIO + +#endif // ELFIO_RELOCATION_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_section.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_section.hpp new file mode 100644 index 0000000000..b2c9b456b5 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_section.hpp @@ -0,0 +1,296 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_SECTION_HPP +#define ELFIO_SECTION_HPP + +#include +#include + +namespace ELFIO { + +class section +{ + friend class elfio; + public: + virtual ~section() {}; + + ELFIO_GET_ACCESS_DECL ( Elf_Half, index ); + ELFIO_GET_SET_ACCESS_DECL( std::string, name ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, type ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, flags ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, info ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, link ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, addr_align ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, entry_size ); + ELFIO_GET_SET_ACCESS_DECL( Elf64_Addr, address ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, size ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, name_string_offset ); + + virtual const char* get_data() const = 0; + virtual void set_data( const char* pData, Elf_Word size ) = 0; + virtual void set_data( const std::string& data ) = 0; + virtual void append_data( const char* pData, Elf_Word size ) = 0; + virtual void append_data( const std::string& data ) = 0; + + protected: + ELFIO_GET_SET_ACCESS_DECL( Elf64_Off, offset ); + ELFIO_SET_ACCESS_DECL( Elf_Half, index ); + + virtual void load( std::istream& f, + std::streampos header_offset ) = 0; + virtual void save( std::ostream& f, + std::streampos header_offset, + std::streampos data_offset ) = 0; + virtual bool is_address_initialized() const = 0; +}; + + +template< class T > +class section_impl : public section +{ + public: +//------------------------------------------------------------------------------ + section_impl( const endianess_convertor* convertor_ ) : convertor( convertor_ ) + { + std::fill_n( reinterpret_cast( &header ), sizeof( header ), '\0' ); + is_address_set = false; + data = 0; + data_size = 0; + } + +//------------------------------------------------------------------------------ + ~section_impl() + { + delete [] data; + } + +//------------------------------------------------------------------------------ + // Section info functions + ELFIO_GET_SET_ACCESS( Elf_Word, type, header.sh_type ); + ELFIO_GET_SET_ACCESS( Elf_Xword, flags, header.sh_flags ); + ELFIO_GET_SET_ACCESS( Elf_Xword, size, header.sh_size ); + ELFIO_GET_SET_ACCESS( Elf_Word, link, header.sh_link ); + ELFIO_GET_SET_ACCESS( Elf_Word, info, header.sh_info ); + ELFIO_GET_SET_ACCESS( Elf_Xword, addr_align, header.sh_addralign ); + ELFIO_GET_SET_ACCESS( Elf_Xword, entry_size, header.sh_entsize ); + ELFIO_GET_SET_ACCESS( Elf_Word, name_string_offset, header.sh_name ); + ELFIO_GET_ACCESS ( Elf64_Addr, address, header.sh_addr ); + +//------------------------------------------------------------------------------ + Elf_Half + get_index() const + { + return index; + } + + +//------------------------------------------------------------------------------ + std::string + get_name() const + { + return name; + } + +//------------------------------------------------------------------------------ + void + set_name( std::string name_ ) + { + name = name_; + } + +//------------------------------------------------------------------------------ + void + set_address( Elf64_Addr value ) + { + header.sh_addr = value; + header.sh_addr = (*convertor)( header.sh_addr ); + is_address_set = true; + } + +//------------------------------------------------------------------------------ + bool + is_address_initialized() const + { + return is_address_set; + } + +//------------------------------------------------------------------------------ + const char* + get_data() const + { + return data; + } + +//------------------------------------------------------------------------------ + void + set_data( const char* raw_data, Elf_Word size ) + { + if ( get_type() != SHT_NOBITS ) { + delete [] data; + try { + data = new char[size]; + } catch (const std::bad_alloc&) { + data = 0; + data_size = 0; + size = 0; + } + if ( 0 != data && 0 != raw_data ) { + data_size = size; + std::copy( raw_data, raw_data + size, data ); + } + } + + set_size( size ); + } + +//------------------------------------------------------------------------------ + void + set_data( const std::string& str_data ) + { + return set_data( str_data.c_str(), (Elf_Word)str_data.size() ); + } + +//------------------------------------------------------------------------------ + void + append_data( const char* raw_data, Elf_Word size ) + { + if ( get_type() != SHT_NOBITS ) { + if ( get_size() + size < data_size ) { + std::copy( raw_data, raw_data + size, data + get_size() ); + } + else { + data_size = 2*( data_size + size); + char* new_data; + try { + new_data = new char[data_size]; + } catch (const std::bad_alloc&) { + new_data = 0; + size = 0; + } + if ( 0 != new_data ) { + std::copy( data, data + get_size(), new_data ); + std::copy( raw_data, raw_data + size, new_data + get_size() ); + delete [] data; + data = new_data; + } + } + set_size( get_size() + size ); + } + } + +//------------------------------------------------------------------------------ + void + append_data( const std::string& str_data ) + { + return append_data( str_data.c_str(), (Elf_Word)str_data.size() ); + } + +//------------------------------------------------------------------------------ + protected: +//------------------------------------------------------------------------------ + ELFIO_GET_SET_ACCESS( Elf64_Off, offset, header.sh_offset ); + +//------------------------------------------------------------------------------ + void + set_index( Elf_Half value ) + { + index = value; + } + +//------------------------------------------------------------------------------ + void + load( std::istream& stream, + std::streampos header_offset ) + { + std::fill_n( reinterpret_cast( &header ), sizeof( header ), '\0' ); + stream.seekg( header_offset ); + stream.read( reinterpret_cast( &header ), sizeof( header ) ); + + Elf_Xword size = get_size(); + if ( 0 == data && SHT_NULL != get_type() && SHT_NOBITS != get_type() ) { + try { + data = new char[size]; + } catch (const std::bad_alloc&) { + data = 0; + data_size = 0; + } + if ( 0 != size ) { + stream.seekg( (*convertor)( header.sh_offset ) ); + stream.read( data, size ); + data_size = size; + } + } + } + +//------------------------------------------------------------------------------ + void + save( std::ostream& f, + std::streampos header_offset, + std::streampos data_offset ) + { + if ( 0 != get_index() ) { + header.sh_offset = data_offset; + header.sh_offset = (*convertor)( header.sh_offset ); + } + + save_header( f, header_offset ); + if ( get_type() != SHT_NOBITS && get_type() != SHT_NULL && + get_size() != 0 && data != 0 ) { + save_data( f, data_offset ); + } + } + +//------------------------------------------------------------------------------ + private: +//------------------------------------------------------------------------------ + void + save_header( std::ostream& f, + std::streampos header_offset ) const + { + f.seekp( header_offset ); + f.write( reinterpret_cast( &header ), sizeof( header ) ); + } + +//------------------------------------------------------------------------------ + void + save_data( std::ostream& f, + std::streampos data_offset ) const + { + f.seekp( data_offset ); + f.write( get_data(), get_size() ); + } + +//------------------------------------------------------------------------------ + private: + T header; + Elf_Half index; + std::string name; + char* data; + Elf_Word data_size; + const endianess_convertor* convertor; + bool is_address_set; +}; + +} // namespace ELFIO + +#endif // ELFIO_SECTION_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_segment.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_segment.hpp new file mode 100644 index 0000000000..35f17e939b --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_segment.hpp @@ -0,0 +1,220 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_SEGMENT_HPP +#define ELFIO_SEGMENT_HPP + +#include +#include + +namespace ELFIO { + +class segment +{ + friend class elfio; + public: + virtual ~segment() {}; + + ELFIO_GET_ACCESS_DECL ( Elf_Half, index ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, type ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Word, flags ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, align ); + ELFIO_GET_SET_ACCESS_DECL( Elf64_Addr, virtual_address ); + ELFIO_GET_SET_ACCESS_DECL( Elf64_Addr, physical_address ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, file_size ); + ELFIO_GET_SET_ACCESS_DECL( Elf_Xword, memory_size ); + ELFIO_GET_ACCESS_DECL( Elf64_Off, offset ); + + virtual const char* get_data() const = 0; + + virtual Elf_Half add_section_index( Elf_Half index, Elf_Xword addr_align ) = 0; + virtual Elf_Half get_sections_num() const = 0; + virtual Elf_Half get_section_index_at( Elf_Half num ) const = 0; + virtual bool is_offset_initialized() const = 0; + + protected: + ELFIO_SET_ACCESS_DECL( Elf64_Off, offset ); + ELFIO_SET_ACCESS_DECL( Elf_Half, index ); + + virtual const std::vector& get_sections() const = 0; + virtual void load( std::istream& stream, std::streampos header_offset ) = 0; + virtual void save( std::ostream& f, std::streampos header_offset, + std::streampos data_offset ) = 0; +}; + + +//------------------------------------------------------------------------------ +template< class T > +class segment_impl : public segment +{ + public: +//------------------------------------------------------------------------------ + segment_impl( endianess_convertor* convertor_ ) : + convertor( convertor_ ) + { + is_offset_set = false; + std::fill_n( reinterpret_cast( &ph ), sizeof( ph ), '\0' ); + data = 0; + } + +//------------------------------------------------------------------------------ + virtual ~segment_impl() + { + delete [] data; + } + +//------------------------------------------------------------------------------ + // Section info functions + ELFIO_GET_SET_ACCESS( Elf_Word, type, ph.p_type ); + ELFIO_GET_SET_ACCESS( Elf_Word, flags, ph.p_flags ); + ELFIO_GET_SET_ACCESS( Elf_Xword, align, ph.p_align ); + ELFIO_GET_SET_ACCESS( Elf64_Addr, virtual_address, ph.p_vaddr ); + ELFIO_GET_SET_ACCESS( Elf64_Addr, physical_address, ph.p_paddr ); + ELFIO_GET_SET_ACCESS( Elf_Xword, file_size, ph.p_filesz ); + ELFIO_GET_SET_ACCESS( Elf_Xword, memory_size, ph.p_memsz ); + ELFIO_GET_ACCESS( Elf64_Off, offset, ph.p_offset ); + +//------------------------------------------------------------------------------ + Elf_Half + get_index() const + { + return index; + } + +//------------------------------------------------------------------------------ + const char* + get_data() const + { + return data; + } + +//------------------------------------------------------------------------------ + Elf_Half + add_section_index( Elf_Half sec_index, Elf_Xword addr_align ) + { + sections.push_back( sec_index ); + if ( addr_align > get_align() ) { + set_align( addr_align ); + } + + return (Elf_Half)sections.size(); + } + +//------------------------------------------------------------------------------ + Elf_Half + get_sections_num() const + { + return (Elf_Half)sections.size(); + } + +//------------------------------------------------------------------------------ + Elf_Half + get_section_index_at( Elf_Half num ) const + { + if ( num < sections.size() ) { + return sections[num]; + } + + return -1; + } + +//------------------------------------------------------------------------------ + protected: +//------------------------------------------------------------------------------ + +//------------------------------------------------------------------------------ + void + set_offset( Elf64_Off value ) + { + ph.p_offset = value; + ph.p_offset = (*convertor)( ph.p_offset ); + is_offset_set = true; + } + +//------------------------------------------------------------------------------ + bool + is_offset_initialized() const + { + return is_offset_set; + } + +//------------------------------------------------------------------------------ + const std::vector& + get_sections() const + { + return sections; + } + +//------------------------------------------------------------------------------ + void + set_index( Elf_Half value ) + { + index = value; + } + +//------------------------------------------------------------------------------ + void + load( std::istream& stream, + std::streampos header_offset ) + { + stream.seekg( header_offset ); + stream.read( reinterpret_cast( &ph ), sizeof( ph ) ); + is_offset_set = true; + + if ( PT_NULL != get_type() && 0 != get_file_size() ) { + stream.seekg( (*convertor)( ph.p_offset ) ); + Elf_Xword size = get_file_size(); + try { + data = new char[size]; + } catch (const std::bad_alloc&) { + data = 0; + } + if ( 0 != data ) { + stream.read( data, size ); + } + } + } + +//------------------------------------------------------------------------------ + void save( std::ostream& f, + std::streampos header_offset, + std::streampos data_offset ) + { + ph.p_offset = data_offset; + ph.p_offset = (*convertor)(ph.p_offset); + f.seekp( header_offset ); + f.write( reinterpret_cast( &ph ), sizeof( ph ) ); + } + +//------------------------------------------------------------------------------ + private: + T ph; + Elf_Half index; + char* data; + std::vector sections; + endianess_convertor* convertor; + bool is_offset_set; +}; + +} // namespace ELFIO + +#endif // ELFIO_SEGMENT_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_strings.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_strings.hpp new file mode 100644 index 0000000000..df952a2145 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_strings.hpp @@ -0,0 +1,96 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_STRINGS_HPP +#define ELFIO_STRINGS_HPP + +#include +#include +#include + +namespace ELFIO { + +//------------------------------------------------------------------------------ +class string_section_accessor +{ + public: +//------------------------------------------------------------------------------ + string_section_accessor( section* section_ ) : + string_section( section_ ) + { + } + + +//------------------------------------------------------------------------------ + const char* + get_string( Elf_Word index ) const + { + if ( string_section ) { + if ( index < string_section->get_size() ) { + const char* data = string_section->get_data(); + if ( 0 != data ) { + return data + index; + } + } + } + + return 0; + } + + +//------------------------------------------------------------------------------ + Elf_Word + add_string( const char* str ) + { + Elf_Word current_position = 0; + + if (string_section) { + // Strings are addeded to the end of the current section data + current_position = (Elf_Word)string_section->get_size(); + + if ( current_position == 0 ) { + char empty_string = '\0'; + string_section->append_data( &empty_string, 1 ); + current_position++; + } + string_section->append_data( str, (Elf_Word)std::strlen( str ) + 1 ); + } + + return current_position; + } + + +//------------------------------------------------------------------------------ + Elf_Word + add_string( const std::string& str ) + { + return add_string( str.c_str() ); + } + +//------------------------------------------------------------------------------ + private: + section* string_section; +}; + +} // namespace ELFIO + +#endif // ELFIO_STRINGS_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_symbols.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_symbols.hpp new file mode 100644 index 0000000000..80e498d8d5 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_symbols.hpp @@ -0,0 +1,278 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_SYMBOLS_HPP +#define ELFIO_SYMBOLS_HPP + +namespace ELFIO { + +//------------------------------------------------------------------------------ +class symbol_section_accessor +{ + public: +//------------------------------------------------------------------------------ + symbol_section_accessor( const elfio& elf_file_, section* symbol_section_ ) : + elf_file( elf_file_ ), + symbol_section( symbol_section_ ) + { + find_hash_section(); + } + +//------------------------------------------------------------------------------ + Elf_Xword + get_symbols_num() const + { + Elf_Xword nRet = 0; + if ( 0 != symbol_section->get_entry_size() ) { + nRet = symbol_section->get_size() / symbol_section->get_entry_size(); + } + + return nRet; + } + +//------------------------------------------------------------------------------ + bool + get_symbol( Elf_Xword index, + std::string& name, + Elf64_Addr& value, + Elf_Xword& size, + unsigned char& bind, + unsigned char& type, + Elf_Half& section_index, + unsigned char& other ) const + { + bool ret = false; + + if ( elf_file.get_class() == ELFCLASS32 ) { + ret = generic_get_symbol( index, name, value, size, bind, + type, section_index, other ); + } + else { + ret = generic_get_symbol( index, name, value, size, bind, + type, section_index, other ); + } + + return ret; + } + +//------------------------------------------------------------------------------ + bool + get_symbol( const std::string& name, + Elf64_Addr& value, + Elf_Xword& size, + unsigned char& bind, + unsigned char& type, + Elf_Half& section_index, + unsigned char& other ) const + { + bool ret = false; + + if ( 0 != get_hash_table_index() ) { + Elf_Word nbucket = *(Elf_Word*)hash_section->get_data(); + Elf_Word nchain = *(Elf_Word*)( hash_section->get_data() + + sizeof( Elf_Word ) ); + Elf_Word val = elf_hash( (const unsigned char*)name.c_str() ); + + Elf_Word y = *(Elf_Word*)( hash_section->get_data() + + ( 2 + val % nbucket ) * sizeof( Elf_Word ) ); + std::string str; + get_symbol( y, str, value, size, bind, type, section_index, other ); + while ( str != name && STN_UNDEF != y && y < nchain ) { + y = *(Elf_Word*)( hash_section->get_data() + + ( 2 + nbucket + y ) * sizeof( Elf_Word ) ); + get_symbol( y, str, value, size, bind, type, section_index, other ); + } + if ( str == name ) { + ret = true; + } + } + + return ret; + } + +//------------------------------------------------------------------------------ + Elf_Word + add_symbol( Elf_Word name, Elf64_Addr value, Elf_Xword size, + unsigned char info, unsigned char other, + Elf_Half shndx ) + { + Elf_Word nRet; + + if ( symbol_section->get_size() == 0 ) { + if ( elf_file.get_class() == ELFCLASS32 ) { + nRet = generic_add_symbol( 0, 0, 0, 0, 0, 0 ); + } + else { + nRet = generic_add_symbol( 0, 0, 0, 0, 0, 0 ); + } + } + + if ( elf_file.get_class() == ELFCLASS32 ) { + nRet = generic_add_symbol( name, value, size, info, other, + shndx ); + } + else { + nRet = generic_add_symbol( name, value, size, info, other, + shndx ); + } + + return nRet; + } + +//------------------------------------------------------------------------------ + Elf_Word + add_symbol( Elf_Word name, Elf64_Addr value, Elf_Xword size, + unsigned char bind, unsigned char type, unsigned char other, + Elf_Half shndx ) + { + return add_symbol( name, value, size, ELF_ST_INFO( bind, type ), other, shndx ); + } + +//------------------------------------------------------------------------------ + Elf_Word + add_symbol( string_section_accessor& pStrWriter, const char* str, + Elf64_Addr value, Elf_Xword size, + unsigned char info, unsigned char other, + Elf_Half shndx ) + { + Elf_Word index = pStrWriter.add_string( str ); + return add_symbol( index, value, size, info, other, shndx ); + } + +//------------------------------------------------------------------------------ + Elf_Word + add_symbol( string_section_accessor& pStrWriter, const char* str, + Elf64_Addr value, Elf_Xword size, + unsigned char bind, unsigned char type, unsigned char other, + Elf_Half shndx ) + { + return add_symbol( pStrWriter, str, value, size, ELF_ST_INFO( bind, type ), other, shndx ); + } + +//------------------------------------------------------------------------------ + private: +//------------------------------------------------------------------------------ + void + find_hash_section() + { + hash_section = 0; + hash_section_index = 0; + Elf_Half nSecNo = elf_file.sections.size(); + for ( Elf_Half i = 0; i < nSecNo && 0 == hash_section_index; ++i ) { + const section* sec = elf_file.sections[i]; + if ( sec->get_link() == symbol_section->get_index() ) { + hash_section = sec; + hash_section_index = i; + } + } + } + +//------------------------------------------------------------------------------ + Elf_Half + get_string_table_index() const + { + return (Elf_Half)symbol_section->get_link(); + } + +//------------------------------------------------------------------------------ + Elf_Half + get_hash_table_index() const + { + return hash_section_index; + } + +//------------------------------------------------------------------------------ + template< class T > + bool + generic_get_symbol( Elf_Xword index, + std::string& name, Elf64_Addr& value, + Elf_Xword& size, + unsigned char& bind, unsigned char& type, + Elf_Half& section_index, + unsigned char& other ) const + { + bool ret = false; + + if ( index < get_symbols_num() ) { + const T* pSym = reinterpret_cast( + symbol_section->get_data() + + index * symbol_section->get_entry_size() ); + + const endianess_convertor& convertor = elf_file.get_convertor(); + + section* string_section = elf_file.sections[get_string_table_index()]; + string_section_accessor str_reader( string_section ); + const char* pStr = str_reader.get_string( convertor( pSym->st_name ) ); + if ( 0 != pStr ) { + name = pStr; + } + value = convertor( pSym->st_value ); + size = convertor( pSym->st_size ); + bind = ELF_ST_BIND( pSym->st_info ); + type = ELF_ST_TYPE( pSym->st_info ); + section_index = convertor( pSym->st_shndx ); + other = pSym->st_other; + + ret = true; + } + + return ret; + } + +//------------------------------------------------------------------------------ + template< class T > + Elf_Word + generic_add_symbol( Elf_Word name, Elf64_Addr value, Elf_Xword size, + unsigned char info, unsigned char other, + Elf_Half shndx ) + { + const endianess_convertor& convertor = elf_file.get_convertor(); + + T entry; + entry.st_name = convertor( name ); + entry.st_value = value; + entry.st_value = convertor( entry.st_value ); + entry.st_size = size; + entry.st_size = convertor( entry.st_size ); + entry.st_info = convertor( info ); + entry.st_other = convertor( other ); + entry.st_shndx = convertor( shndx ); + + symbol_section->append_data( reinterpret_cast( &entry ), + sizeof( entry ) ); + + Elf_Word nRet = symbol_section->get_size() / sizeof( entry ) - 1; + + return nRet; + } + +//------------------------------------------------------------------------------ + private: + const elfio& elf_file; + section* symbol_section; + Elf_Half hash_section_index; + const section* hash_section; +}; + +} // namespace ELFIO + +#endif // ELFIO_SYMBOLS_HPP diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_utils.hpp b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_utils.hpp new file mode 100644 index 0000000000..f8423bd147 --- /dev/null +++ b/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_utils.hpp @@ -0,0 +1,209 @@ +/* +Copyright (C) 2001-2015 by Serge Lamikhov-Center + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef ELFIO_UTILS_HPP +#define ELFIO_UTILS_HPP + +#define ELFIO_GET_ACCESS( TYPE, NAME, FIELD ) \ + TYPE get_##NAME() const \ + { \ + return (*convertor)( FIELD ); \ + } +#define ELFIO_SET_ACCESS( TYPE, NAME, FIELD ) \ + void set_##NAME( TYPE value ) \ + { \ + FIELD = value; \ + FIELD = (*convertor)( FIELD ); \ + } +#define ELFIO_GET_SET_ACCESS( TYPE, NAME, FIELD ) \ + TYPE get_##NAME() const \ + { \ + return (*convertor)( FIELD ); \ + } \ + void set_##NAME( TYPE value ) \ + { \ + FIELD = value; \ + FIELD = (*convertor)( FIELD ); \ + } + +#define ELFIO_GET_ACCESS_DECL( TYPE, NAME ) \ + virtual TYPE get_##NAME() const = 0 + +#define ELFIO_SET_ACCESS_DECL( TYPE, NAME ) \ + virtual void set_##NAME( TYPE value ) = 0 + +#define ELFIO_GET_SET_ACCESS_DECL( TYPE, NAME ) \ + virtual TYPE get_##NAME() const = 0; \ + virtual void set_##NAME( TYPE value ) = 0 + +namespace ELFIO { + +//------------------------------------------------------------------------------ +class endianess_convertor { + public: +//------------------------------------------------------------------------------ + endianess_convertor() + { + need_conversion = false; + } + +//------------------------------------------------------------------------------ + void + setup( unsigned char elf_file_encoding ) + { + need_conversion = ( elf_file_encoding != get_host_encoding() ); + } + +//------------------------------------------------------------------------------ + uint64_t + operator()( uint64_t value ) const + { + if ( !need_conversion ) { + return value; + } + value = + ( ( value & 0x00000000000000FFull ) << 56 ) | + ( ( value & 0x000000000000FF00ull ) << 40 ) | + ( ( value & 0x0000000000FF0000ull ) << 24 ) | + ( ( value & 0x00000000FF000000ull ) << 8 ) | + ( ( value & 0x000000FF00000000ull ) >> 8 ) | + ( ( value & 0x0000FF0000000000ull ) >> 24 ) | + ( ( value & 0x00FF000000000000ull ) >> 40 ) | + ( ( value & 0xFF00000000000000ull ) >> 56 ); + + return value; + } + +//------------------------------------------------------------------------------ + int64_t + operator()( int64_t value ) const + { + if ( !need_conversion ) { + return value; + } + return (int64_t)(*this)( (uint64_t)value ); + } + +//------------------------------------------------------------------------------ + uint32_t + operator()( uint32_t value ) const + { + if ( !need_conversion ) { + return value; + } + value = + ( ( value & 0x000000FF ) << 24 ) | + ( ( value & 0x0000FF00 ) << 8 ) | + ( ( value & 0x00FF0000 ) >> 8 ) | + ( ( value & 0xFF000000 ) >> 24 ); + + return value; + } + +//------------------------------------------------------------------------------ + int32_t + operator()( int32_t value ) const + { + if ( !need_conversion ) { + return value; + } + return (int32_t)(*this)( (uint32_t)value ); + } + +//------------------------------------------------------------------------------ + uint16_t + operator()( uint16_t value ) const + { + if ( !need_conversion ) { + return value; + } + value = + ( ( value & 0x00FF ) << 8 ) | + ( ( value & 0xFF00 ) >> 8 ); + + return value; + } + +//------------------------------------------------------------------------------ + int16_t + operator()( int16_t value ) const + { + if ( !need_conversion ) { + return value; + } + return (int16_t)(*this)( (uint16_t)value ); + } + +//------------------------------------------------------------------------------ + int8_t + operator()( int8_t value ) const + { + return value; + } + +//------------------------------------------------------------------------------ + uint8_t + operator()( uint8_t value ) const + { + return value; + } + +//------------------------------------------------------------------------------ + private: +//------------------------------------------------------------------------------ + unsigned char + get_host_encoding() const + { + static const int tmp = 1; + if ( 1 == *(char*)&tmp ) { + return ELFDATA2LSB; + } + else { + return ELFDATA2MSB; + } + } + +//------------------------------------------------------------------------------ + private: + bool need_conversion; +}; + + +//------------------------------------------------------------------------------ +inline +uint32_t +elf_hash( const unsigned char *name ) +{ + uint32_t h = 0, g; + while ( *name ) { + h = (h << 4) + *name++; + g = h & 0xf0000000; + if ( g != 0 ) + h ^= g >> 24; + h &= ~g; + } + return h; +} + +} // namespace ELFIO + +#endif // ELFIO_UTILS_HPP diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 74b0092b84..e3dc511d02 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -23,15 +23,18 @@ THE SOFTWARE. #include #include #include -#include -#include -#include +#include +#include +#include +#include +#include #include -#include "hsa/hsa.h" -#include "hsa/hsa_ext_amd.h" -#include "hsa/amd_hsa_kernel_code.h" +#include +#include +#include +#include "hip/hcc_detail/elfio/elfio.hpp" #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" #include "trace_helper.h" @@ -114,6 +117,8 @@ namespace hipdrv { } // End namespace hipdrv uint64_t PrintSymbolSizes(const void *emi, const char *name){ + using namespace ELFIO; + const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; if(NULL == ehdr || EV_CURRENT != ehdr->e_version){} const Elf64_Shdr * shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); @@ -138,6 +143,8 @@ uint64_t PrintSymbolSizes(const void *emi, const char *name){ } uint64_t ElfSize(const void *emi){ + using namespace ELFIO; + const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); @@ -157,8 +164,155 @@ uint64_t ElfSize(const void *emi){ return total_size; } +namespace +{ + template + inline + ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) + { + using namespace std; + + const auto it = find_if( + reader.sections.begin(), reader.sections.end(), move(p)); + + return it != reader.sections.end() ? *it : nullptr; + } + + inline + std::vector copy_names_of_undefined_symbols( + const ELFIO::symbol_section_accessor& section) + { + using namespace ELFIO; + using namespace std; + + vector r; + + for (auto i = 0u; i != section.get_symbols_num(); ++i) { + // TODO: this is boyscout code, caching the temporaries + // may be of worth. + string name; + Elf64_Addr value = 0; + Elf_Xword size = 0; + Elf_Half sect_idx = 0; + uint8_t bind = 0; + uint8_t type = 0; + uint8_t other = 0; + + section.get_symbol( + i, name, value, size, bind, type, sect_idx, other); + + if (sect_idx == SHN_UNDEF && !name.empty()) { + r.push_back(std::move(name)); + } + } + + return r; + } + + inline + std::pair find_symbol_address( + const ELFIO::symbol_section_accessor& section, + const std::string& symbol_name) + { + using namespace ELFIO; + using namespace std; + + static constexpr pair r{0, 0}; + + for (auto i = 0u; i != section.get_symbols_num(); ++i) { + // TODO: this is boyscout code, caching the temporaries + // may be of worth. + string name; + Elf64_Addr value = 0; + Elf_Xword size = 0; + Elf_Half sect_idx = 0; + uint8_t bind = 0; + uint8_t type = 0; + uint8_t other = 0; + + section.get_symbol( + i, name, value, size, bind, type, sect_idx, other); + + if (name == symbol_name) return make_pair(value, size); + } + + return r; + } + + inline + void associate_code_object_symbols_with_host_allocation( + const ELFIO::elfio& reader, + ELFIO::section* code_object_dynsym, + ELFIO::section* process_symtab, + hsa_agent_t agent, + hsa_executable_t executable) + { + using namespace ELFIO; + using namespace std; + + if (!code_object_dynsym || !process_symtab) return; + + const auto undefined_symbols = copy_names_of_undefined_symbols( + symbol_section_accessor{reader, code_object_dynsym}); + + for (auto&& x : undefined_symbols) { + const auto tmp = find_symbol_address( + symbol_section_accessor{reader, process_symtab}, x); + + assert(tmp.first); + + void* p = nullptr; + hsa_amd_memory_lock( + reinterpret_cast(tmp.first), tmp.second, &agent, 1, &p); + + hsa_executable_agent_global_variable_define( + executable, agent, x.c_str(), p); + + static vector< + unique_ptr> globals; + static mutex mtx; + + lock_guard lck{mtx}; + globals.emplace_back(p, hsa_amd_memory_unlock); + } + } + + inline + void load_code_object_and_freeze_executable( + const char* file, hsa_agent_t agent, hsa_executable_t executable) + { // TODO: the following sequence is inefficient, should be refactored + // into a single load of the file and subsequent ELFIO + // processing. + using namespace std; + + static const auto cor_deleter = [](hsa_code_object_reader_t* p) { + hsa_code_object_reader_destroy(*p); + }; + + using RAII_code_reader = unique_ptr< + hsa_code_object_reader_t, decltype(cor_deleter)>; + + unique_ptr cobj{fopen(file, "r"), fclose}; + RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter}; + hsa_code_object_reader_create_from_file(fileno(cobj.get()), tmp.get()); + + hsa_executable_load_agent_code_object( + executable, agent, *tmp, nullptr, nullptr); + + hsa_executable_freeze(executable, nullptr); + + static vector code_readers; + static mutex mtx; + + lock_guard lck{mtx}; + code_readers.push_back(move(tmp)); + } +} + +hipError_t hipModuleLoad(hipModule_t *module, const char *fname) +{ + using namespace ELFIO; -hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ HIP_INIT_API(module, fname); hipError_t ret = hipSuccess; *module = new ihipModule_t; @@ -174,96 +328,40 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ }else{ int deviceId = ctx->getDevice()->_deviceId; ihipDevice_t *currentDevice = ihipGetDevice(deviceId); - std::ifstream in(fname, std::ios::binary | std::ios::ate); - if(!in.is_open() ){ + hsa_executable_create_alt( + HSA_PROFILE_FULL, + HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, + nullptr, + &(*module)->executable); + + elfio reader; + if (!reader.load(fname)) { return ihipLogStatus(hipErrorFileNotFound); + } + else { + // TODO: this may benefit from caching as well. + elfio self_reader; + self_reader.load("/proc/self/exe"); + const auto symtab = + find_section_if(self_reader, [](const ELFIO::section* x) { + return x->get_type() == SHT_SYMTAB; + }); - } else { + const auto code_object_dynsym = + find_section_if(reader, [](const ELFIO::section* x) { + return x->get_type() == SHT_DYNSYM; + }); - size_t size = std::string::size_type(in.tellg()); - void *p = NULL; - hsa_agent_t agent = currentDevice->_hsaAgent; - hsa_region_t sysRegion; - hsa_status_t status = hsa_agent_iterate_regions(agent, hipdrv::findSystemRegions, &sysRegion); - status = hsa_memory_allocate(sysRegion, size, (void**)&p); + associate_code_object_symbols_with_host_allocation( + reader, + code_object_dynsym, + symtab, + currentDevice->_hsaAgent, + (*module)->executable); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorOutOfMemory); - } - - char *ptr = (char*)p; - if(!ptr){ - return ihipLogStatus(hipErrorOutOfMemory); - } - (*module)->ptr = p; - (*module)->size = size; - in.seekg(0, std::ios::beg); - std::copy(std::istreambuf_iterator(in), std::istreambuf_iterator(), ptr); -/* - Elf *e = elf_memory((char*)p, size); - if(elf_kind(e) != ELF_K_ELF){ - return ihipLogStatus(hipErrorInvalidValue); - } - size_t numpHdrs; - if(elf_getphdrnum(e, &numpHdrs) != 0){ - return ihipLogStatus(hipErrorInvalidValue); - } - for(size_t i=0;i= sizeof(int)){ - char* ptr = (char*) p + pHdr.p_offset; - char* segmentEnd = ptr + pHdr.p_filesz; - while(ptr < segmentEnd){ - MyElfNote *note = (MyElfNote*) ptr; - char *name = (char*) ¬e[1]; - char *desc = name + alignTo(note->n_namesz, sizeof(int)); - if (note->n_type == 8) { - std::string metadatastr((const char*)desc, (size_t)note->n_descsz); - AMDGPU::RuntimeMD::Program::Metadata meta; - meta = meta.fromYAML(metadatastr); - for(int i=0;in_namesz, sizeof(int)) - + alignTo(note->n_descsz, sizeof(int)); - } - } - } -*/ - status = hsa_code_object_deserialize(ptr, size, NULL, &(*module)->object); - - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorSharedObjectInitFailed); - } - - status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - CHECKLOG_HSA(status, hipErrorNotInitialized); - - status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); - CHECKLOG_HSA(status, hipErrorNotInitialized); - - status = hsa_executable_freeze((*module)->executable, NULL); - CHECKLOG_HSA(status, hipErrorNotInitialized); + load_code_object_and_freeze_executable( + fname, currentDevice->_hsaAgent, (*module)->executable); } } @@ -283,16 +381,16 @@ hipError_t hipModuleUnload(hipModule_t hmod) { ret = hipErrorInvalidValue; } - status = hsa_code_object_destroy(hmod->object); - if(status != HSA_STATUS_SUCCESS) - { - ret = hipErrorInvalidValue; - } - status = hsa_memory_free(hmod->ptr); - if(status != HSA_STATUS_SUCCESS) - { - ret = hipErrorInvalidValue; - } + // status = hsa_code_object_destroy(hmod->object); + // if(status != HSA_STATUS_SUCCESS) + // { + // ret = hipErrorInvalidValue; + // } + // status = hsa_memory_free(hmod->ptr); + // if(status != HSA_STATUS_SUCCESS) + // { + // ret = hipErrorInvalidValue; + // } for(auto f = hmod->funcTrack.begin(); f != hmod->funcTrack.end(); ++f) { delete *f; } @@ -404,7 +502,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, } else { return ihipLogStatus(hipErrorNotInitialized); } - + }else{ return ihipLogStatus(hipErrorInvalidValue); } @@ -450,7 +548,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, hc::completion_future cf; - lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, + lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize, (startEvent || stopEvent) ? &cf : nullptr #if (__hcc_workweek__ > 17312) , f->_name.c_str() From 1c8db33de94d4a164b6872ae00ba45f163bddab6 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 9 Oct 2017 15:46:38 +0100 Subject: [PATCH 02/46] This fixes incorrect usage of the reader object, which created arcane mismatches when one reader accessed another's section. [ROCm/clr commit: b222448efd1253fa5ea184dfb063363c815cf189] --- projects/clr/hipamd/src/hip_module.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index e3dc511d02..b6bc9d63fd 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -242,6 +242,7 @@ namespace inline void associate_code_object_symbols_with_host_allocation( const ELFIO::elfio& reader, + const ELFIO::elfio& self_reader, ELFIO::section* code_object_dynsym, ELFIO::section* process_symtab, hsa_agent_t agent, @@ -257,7 +258,7 @@ namespace for (auto&& x : undefined_symbols) { const auto tmp = find_symbol_address( - symbol_section_accessor{reader, process_symtab}, x); + symbol_section_accessor{self_reader, process_symtab}, x); assert(tmp.first); @@ -343,6 +344,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname) // TODO: this may benefit from caching as well. elfio self_reader; self_reader.load("/proc/self/exe"); + const auto symtab = find_section_if(self_reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; @@ -355,6 +357,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname) associate_code_object_symbols_with_host_allocation( reader, + self_reader, code_object_dynsym, symtab, currentDevice->_hsaAgent, From c3f743a7631fcf8e76ea648de9ec9362f8643a19 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Thu, 19 Oct 2017 13:34:02 +0100 Subject: [PATCH 03/46] Add [[nodiscard]] attribute to hipError_t in C++17 mode This will yield a compiler warning any time you discard an error return value from an API call. [ROCm/clr commit: 939fa02ade27cbde15bd6fe4f999952afd00b9ac] --- projects/clr/hipamd/include/hip/hip_runtime_api.h | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index 97733964bf..a2ebcd27a5 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -143,6 +143,13 @@ typedef struct hipPointerAttribute_t { * */ +// Ignoring error-code return values from hip APIs is discouraged. On C++17, +// we can make that yield a warning +#if __cplusplus >= 201703L +#define __HIP_NODISCARD [[nodiscard]] +#else +#define __HIP_NODISCARD +#endif /* * @brief hipError_t @@ -152,7 +159,7 @@ typedef struct hipPointerAttribute_t { // Developer note - when updating these, update the hipErrorName and hipErrorString functions in NVCC and HCC paths // Also update the hipCUDAErrorTohipError function in NVCC path. -typedef enum hipError_t { +typedef enum __HIP_NODISCARD hipError_t { hipSuccess = 0, ///< Successful completion. hipErrorOutOfMemory = 2, hipErrorNotInitialized = 3, @@ -219,6 +226,8 @@ typedef enum hipError_t { hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; +#undef __HIP_NODISCARD + /* * @brief hipDeviceAttribute_t * @enum From a0c26d39f93ae4f4dc450fdcee07f4e189c6e020 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:16:11 +0100 Subject: [PATCH 04/46] Fix broken indentation introduced by previous commit [ROCm/clr commit: 72c9767b1dd61d4eef569ba4f866ea23ab82c8b9] --- .../clr/hipamd/hipify-clang/CMakeLists.txt | 82 +++++++++---------- 1 file changed, 41 insertions(+), 41 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 1aaa68097f..8bb366a661 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -13,6 +13,7 @@ if (NOT ${LLVM_FOUND}) message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") endif() endif() + if (${LLVM_FOUND}) list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) include(AddLLVM) @@ -52,55 +53,54 @@ if (${LLVM_FOUND}) LLVMOption LLVMCore) -if(WIN32) - target_link_libraries(hipify-clang version) -endif() + if(WIN32) + target_link_libraries(hipify-clang version) + endif() set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") -if(MSVC) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") - set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") -else() - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") -endif() + if(MSVC) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") + set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") + else() + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") + endif() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") install(TARGETS hipify-clang DESTINATION bin) -if (HIPIFY_CLANG_TESTS) - # tests - set(Python_ADDITIONAL_VERSIONS 2.7) - include(FindPythonInterp) - if(NOT PYTHONINTERP_FOUND) - message(FATAL_ERROR - "Unable to find Python interpreter, required for builds and testing.\n\n" - "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.") + if (HIPIFY_CLANG_TESTS) + # tests + set(Python_ADDITIONAL_VERSIONS 2.7) + include(FindPythonInterp) + if(NOT PYTHONINTERP_FOUND) + message(FATAL_ERROR + "Unable to find Python interpreter, required for builds and testing.\n\n" + "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.") + endif() + + find_program(LIT_COMMAND lit) + + set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + + configure_file( + ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + @ONLY) + + add_lit_testsuite(test-hipify "Running HIPify regression tests" + ${CMAKE_SOURCE_DIR}/tests/hipify-clang + PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + DEPENDS hipify-clang lit + ) + + add_custom_target(test-hipify-clang) + add_dependencies(test-hipify-clang test-hipify) + set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") endif() - find_program(LIT_COMMAND lit) - - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) - - configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) - - add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit - ) - - add_custom_target(test-hipify-clang) - add_dependencies(test-hipify-clang test-hipify) - set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") -endif() - -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) -endif() - + if (PARENT_SCOPE) + set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) + endif() endif() From 28ebdf5379957a3531d10991835484a33823898c Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:17:29 +0100 Subject: [PATCH 05/46] Use early return to avoid indenting all of CMakeLists.txt [ROCm/clr commit: d3a56460512bfcbf5289ceb08ede50355aa633df] --- .../clr/hipamd/hipify-clang/CMakeLists.txt | 167 +++++++++--------- 1 file changed, 83 insertions(+), 84 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 8bb366a661..7cfac3301a 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -11,96 +11,95 @@ if (NOT ${LLVM_FOUND}) find_package(LLVM 3.9 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) if (NOT ${LLVM_FOUND}) message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") + return() endif() endif() -if (${LLVM_FOUND}) - list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) - include(AddLLVM) +list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) +include(AddLLVM) - message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") +message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") - include_directories(${LLVM_INCLUDE_DIRS}) - link_directories(${LLVM_LIBRARY_DIRS}) - add_definitions(${LLVM_DEFINITIONS}) - add_llvm_executable(hipify-clang src/Cuda2Hip.cpp) +include_directories(${LLVM_INCLUDE_DIRS}) +link_directories(${LLVM_LIBRARY_DIRS}) +add_definitions(${LLVM_DEFINITIONS}) +add_llvm_executable(hipify-clang src/Cuda2Hip.cpp) - set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) - set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) +set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) +set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) - # Link against LLVM and CLANG libraries - target_link_libraries(hipify-clang - clangASTMatchers - clangFrontend - clangTooling - clangParse - clangSerialization - clangSema - clangEdit - clangFormat - clangLex - clangAnalysis - clangDriver - clangAST - clangToolingCore - clangRewrite - clangBasic - LLVMProfileData - LLVMSupport - LLVMMCParser - LLVMMC - LLVMBitReader - LLVMOption - LLVMCore) +# Link against LLVM and CLANG libraries +target_link_libraries(hipify-clang + clangASTMatchers + clangFrontend + clangTooling + clangParse + clangSerialization + clangSema + clangEdit + clangFormat + clangLex + clangAnalysis + clangDriver + clangAST + clangToolingCore + clangRewrite + clangBasic + LLVMProfileData + LLVMSupport + LLVMMCParser + LLVMMC + LLVMBitReader + LLVMOption + LLVMCore) - if(WIN32) - target_link_libraries(hipify-clang version) - endif() - - set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") - if(MSVC) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") - set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") - else() - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") - endif() - - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") - - install(TARGETS hipify-clang DESTINATION bin) - - if (HIPIFY_CLANG_TESTS) - # tests - set(Python_ADDITIONAL_VERSIONS 2.7) - include(FindPythonInterp) - if(NOT PYTHONINTERP_FOUND) - message(FATAL_ERROR - "Unable to find Python interpreter, required for builds and testing.\n\n" - "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.") - endif() - - find_program(LIT_COMMAND lit) - - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) - - configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) - - add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit - ) - - add_custom_target(test-hipify-clang) - add_dependencies(test-hipify-clang test-hipify) - set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") - endif() - - if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) - endif() +if(WIN32) + target_link_libraries(hipify-clang version) +endif() + +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") +if(MSVC) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") + set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") +else() + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") +endif() + +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") + +install(TARGETS hipify-clang DESTINATION bin) + +if (HIPIFY_CLANG_TESTS) + # tests + set(Python_ADDITIONAL_VERSIONS 2.7) + include(FindPythonInterp) + if(NOT PYTHONINTERP_FOUND) + message(FATAL_ERROR + "Unable to find Python interpreter, required for builds and testing.\n\n" + "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.") + endif() + + find_program(LIT_COMMAND lit) + + set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + + configure_file( + ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + @ONLY) + + add_lit_testsuite(test-hipify "Running HIPify regression tests" + ${CMAKE_SOURCE_DIR}/tests/hipify-clang + PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + DEPENDS hipify-clang lit + ) + + add_custom_target(test-hipify-clang) + add_dependencies(test-hipify-clang test-hipify) + set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") +endif() + +if (PARENT_SCOPE) + set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) endif() From 9eac4a20123e6c5ad9c052ef841096406520d3d4 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:18:06 +0100 Subject: [PATCH 06/46] Declare HIPIFY_CLANG_TESTS as a cmake option [ROCm/clr commit: b926cdf0028ea67c50cb9a9b44200ba5ae92a69a] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 7cfac3301a..21efb5bc80 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -1,6 +1,8 @@ cmake_minimum_required(VERSION 2.8.8) project(hipify-clang) +option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) + if (PARENT_SCOPE) set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) endif() From 35943634ea28642da0be59340f5c62d9ba7766f0 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:21:47 +0100 Subject: [PATCH 07/46] Don't be picky about clang versions [ROCm/clr commit: 03dd8c9fcad558f6b20c661de59ae9083d10eec2] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 10 +--------- 1 file changed, 1 insertion(+), 9 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 21efb5bc80..6e3f4e80de 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -7,15 +7,7 @@ if (PARENT_SCOPE) set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) endif() -# Find LLVM package -find_package(LLVM 3.8 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) -if (NOT ${LLVM_FOUND}) - find_package(LLVM 3.9 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) - if (NOT ${LLVM_FOUND}) - message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") - return() - endif() -endif() +find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) include(AddLLVM) From 95e170653799f05f84042e3b411ea1ae0bea45ad Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:22:12 +0100 Subject: [PATCH 08/46] Don't reinvent find_package PythonInterp is a finder module that ships with cmake. It supports the conventional interaction with find_package that allows you to demand success, and particular vesions, without having your own logic: https://cmake.org/cmake/help/v3.0/command/find_package.html [ROCm/clr commit: c62766f880eae1d423c08bacc43e7c9d9c18bca1] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 6e3f4e80de..4c6dc8aec8 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -66,14 +66,7 @@ install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) # tests - set(Python_ADDITIONAL_VERSIONS 2.7) - include(FindPythonInterp) - if(NOT PYTHONINTERP_FOUND) - message(FATAL_ERROR - "Unable to find Python interpreter, required for builds and testing.\n\n" - "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.") - endif() - + find_package(PythonInterp 2.7 REQUIRED EXACT) find_program(LIT_COMMAND lit) set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) From f736f861b8b0a930b243bbeac8587b4f944cc4ea Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:23:32 +0100 Subject: [PATCH 09/46] Skip `lit` tests if `lit` cannot be found [ROCm/clr commit: 6abd8963aed63fb89217da086cfefb99f559eed7] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 4c6dc8aec8..718e3d15de 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -3,6 +3,13 @@ project(hipify-clang) option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) +# Disable the tests if `lit` is not installed. +find_program(LIT_COMMAND lit) +if (NOT LIT_COMMAND) + set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") + message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") +endif() + if (PARENT_SCOPE) set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) endif() @@ -67,20 +74,20 @@ install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) # tests find_package(PythonInterp 2.7 REQUIRED EXACT) - find_program(LIT_COMMAND lit) set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) configure_file( ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) + @ONLY + ) add_lit_testsuite(test-hipify "Running HIPify regression tests" ${CMAKE_SOURCE_DIR}/tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg DEPENDS hipify-clang lit - ) + ) add_custom_target(test-hipify-clang) add_dependencies(test-hipify-clang test-hipify) From 37764de42834e1eb1edc1ff7c44ec8836fc23af9 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:24:35 +0100 Subject: [PATCH 10/46] Use the cache for global variables - not PARENT_SCOPE hacks [ROCm/clr commit: ed6681f5635dd61def669c47c22dfa8ec0c07c56] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 718e3d15de..355b0cea7a 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -10,9 +10,7 @@ if (NOT LIT_COMMAND) message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") endif() -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) -endif() +set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) @@ -94,6 +92,4 @@ if (HIPIFY_CLANG_TESTS) set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") endif() -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) -endif() +set(BUILD_HIPIFY_CLANG 1 CACHE INTERNAL "") From 17903e9dbd1c35f053ff4869a5e2a3f53c9f20c7 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:09:48 +0100 Subject: [PATCH 11/46] Fix two faulty LLVM version checks What we actually want to do here is use the StringRef version in versions newer than 3.8, and the void one in 3.8 and older. Checking "major-version >= 3 && minor-version >= 9" does not do what we want. Consider what this will do for version 4.0, for which minor-version is zero... [ROCm/clr commit: c876f6ffd5379c5c3bc24257db99ac392390224f] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 907e78e54f..5e03e9732d 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3109,10 +3109,10 @@ public: // to workaround the 'const' MacroArgs passed into this hook. const Token *start = Args->getUnexpArgument(i); size_t len = Args->getArgLength(start) + 1; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - _pp->EnterTokenStream(ArrayRef(start, len), false); -#else +#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) _pp->EnterTokenStream(start, len, false, false); +#else + _pp->EnterTokenStream(ArrayRef(start, len), false); #endif do { toks.push_back(Token()); @@ -4207,11 +4207,15 @@ void copyFile(const std::string& src, const std::string& dst) { int main(int argc, const char **argv) { auto start = std::chrono::steady_clock::now(); auto begin = start; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - llvm::sys::PrintStackTraceOnErrorSignal(StringRef()); -#else + + // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support + // anything older than 3.8, so let's specifically detect the one old version we support. +#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) llvm::sys::PrintStackTraceOnErrorSignal(); +#else + llvm::sys::PrintStackTraceOnErrorSignal(StringRef()); #endif + CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore); std::vector fileSources = OptionsParser.getSourcePathList(); std::string dst = OutputFilename; From 3f284c0927e23b3d2327cde09fab19e5b79660ef Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:16:39 +0100 Subject: [PATCH 12/46] Refer to clang::StringLiteral explicitly Newer versions of llvm/clang mean there is both an llvm::StringLiteral and a clang::StringLiteral. Since we're dumping both namespaces wholesale into the global namespace with `using` declarations, this creates a name collision, which must be resolved. This change is backwards-compatible, and fixes a problem you encounter when using newer versions of the llvm/clang API. [ROCm/clr commit: 73984ed8094db75070a37c7120cf7c4cc8b57081] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 5e03e9732d..cf11dc9546 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3799,7 +3799,7 @@ private: } bool stringLiteral(const MatchFinder::MatchResult &Result) { - if (const StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { + if (const clang::StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { if (sLiteral->getCharByteWidth() == 1) { StringRef s = sLiteral->getString(); SourceManager *SM = Result.SourceManager; From ae40d4bdea59a2af6f09429ab666f5870776654d Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:29:47 +0100 Subject: [PATCH 13/46] Remove unnecessary call to Retain() on a smart pointer The Preprocessor smart pointer is held by the CompilerInstance, and therefore its reference count cannot reach zero until the CompilerInstance itself is destroyed. If the CompilerInstance is destroyed, you have more to worry about than just the preprocessor being deallocated! Newer versions of the LLVM/Clang API migrated to using std::shared_ptr, so there is no `Retain()` function (by that name, anyway). Eliminating this redundant use is a neat and backward-compatible way to become compatible with newer versions of the LLVM/Clang API. [ROCm/clr commit: 22a5e2330d94257cd45c8b611b641aab71427ea5] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index cf11dc9546..fc224adaf0 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3029,7 +3029,6 @@ public: SourceManager &SM = CI.getSourceManager(); setSourceManager(&SM); PP.addPPCallbacks(std::unique_ptr(this)); - PP.Retain(); setPreprocessor(&PP); return true; } From db2b03b879f8ed803e2ee3270a121474ac70e260 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:38:24 +0100 Subject: [PATCH 14/46] Cope with clang 4.0's rename of getNumArgs() Sorry, this one I couldn't do in a perfectly elegant way ;) [ROCm/clr commit: 9ccb966e2bc02c1745fd8d1b521efdc5490ebfe1] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index fc224adaf0..ef53afaaf3 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3102,7 +3102,14 @@ public: const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { - for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs(); i++) { + // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics + // remain unchanged. +#if LLVM_VERSION_MAJOR > 3 + #define GET_NUM_ARGS() getNumParams() +#else + #define GET_NUM_ARGS() getNumArgs() +#endif + for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) { std::vector toks; // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument' // to workaround the 'const' MacroArgs passed into this hook. From a21737be3964e1cd22c9c927f21143e3108200b2 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:56:31 +0100 Subject: [PATCH 15/46] Cope with Replacements now having llvm::Error returns [ROCm/clr commit: 885573e1e2fba3db76d8a7d101eb6a16806dcc4c] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index ef53afaaf3..816ca4bc3d 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -2928,7 +2928,13 @@ protected: std::string mainFileName; virtual void insertReplacement(const Replacement &rep, const FullSourceLoc &fullSL) { +#if LLVM_VERSION_MAJOR > 3 + // New clang added error checking to Replacements, and *insists* that you explicitly check it. + llvm::Error e = Replace->add(rep); +#else + // In older versions, it's literally an std::set Replace->insert(rep); +#endif if (PrintStats) { LOCs.insert(fullSL.getExpansionLineNumber()); } From fd9bc0d0d82ef717e55a4faf12b3e5707bfb74b8 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:56:46 +0100 Subject: [PATCH 16/46] Omit now-removed Filename string arg from handleBeginSource [ROCm/clr commit: abd34b2125403f62d5778a9446445eb10304cd67] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 816ca4bc3d..0f095c7245 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3030,7 +3030,11 @@ public: HipifyPPCallbacks(Replacements *R, const std::string &mainFileName) : Cuda2Hip(R, mainFileName), SeenEnd(false), _sm(nullptr), _pp(nullptr) {} - virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override { + virtual bool handleBeginSource(CompilerInstance &CI +#if LLVM_VERSION_MAJOR < 4 + , StringRef Filename +#endif + ) override { Preprocessor &PP = CI.getPreprocessor(); SourceManager &SM = CI.getSourceManager(); setSourceManager(&SM); From 5e09ca10b78f78596559603bb16b0d87f186662d Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 11:09:17 +0100 Subject: [PATCH 17/46] Use inline initialisers to set default field values A trivial cleanup that helps in a moment.. [ROCm/clr commit: 9d83348ac72addf50a3487fe2c33a1a5d9d37322] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 0f095c7245..b71a55d81d 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3028,7 +3028,7 @@ class Cuda2HipCallback; class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip { public: HipifyPPCallbacks(Replacements *R, const std::string &mainFileName) - : Cuda2Hip(R, mainFileName), SeenEnd(false), _sm(nullptr), _pp(nullptr) {} + : Cuda2Hip(R, mainFileName) {} virtual bool handleBeginSource(CompilerInstance &CI #if LLVM_VERSION_MAJOR < 4 @@ -3201,15 +3201,15 @@ public: void EndOfMainFile() override {} - bool SeenEnd; + bool SeenEnd = false; void setSourceManager(SourceManager *sm) { _sm = sm; } void setPreprocessor(Preprocessor *pp) { _pp = pp; } void setMatch(Cuda2HipCallback *match) { Match = match; } private: - SourceManager *_sm; - Preprocessor *_pp; - Cuda2HipCallback *Match; + SourceManager *_sm = nullptr; + Preprocessor *_pp = nullptr; + Cuda2HipCallback *Match = nullptr; }; class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip { From 64a26027b79587ad562286966220c4d23586449b Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 11:43:35 +0100 Subject: [PATCH 18/46] Be agnostic to the new getReplacements() API See comment [ROCm/clr commit: a06fe54ee8027a6530424c6841b70618846b0bd3] --- .../clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 26 +++++++++++++++---- 1 file changed, 21 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index b71a55d81d..8b16e75db0 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -4291,8 +4291,19 @@ int main(int argc, const char **argv) { // So what we do is operate on a copy, which we then move to the output. RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile); ast_matchers::MatchFinder Finder; - HipifyPPCallbacks PPCallbacks(&Tool.getReplacements(), tmpFile); - Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks, tmpFile); + + // The Replacements to apply to the file `src`. + Replacements* replacementsToUse; +#if LLVM_VERSION_MAJOR > 3 + // getReplacements() now returns a map from filename to Replacements - so create an entry + // for this source file and return a pointer to it. + replacementsToUse = &(Tool.getReplacements()[tmpFile]); +#else + replacementsToUse = &Tool.getReplacements(); +#endif + + HipifyPPCallbacks PPCallbacks(replacementsToUse, tmpFile); + Cuda2HipCallback Callback(replacementsToUse, &Finder, &PPCallbacks, tmpFile); addAllMatchers(Finder, &Callback); @@ -4320,9 +4331,14 @@ int main(int argc, const char **argv) { SourceManager SM(Diagnostics, Tool.getFiles()); if (PrintStats) { DEBUG(dbgs() << "Replacements collected by the tool:\n"); - for (const auto &r : Tool.getReplacements()) { - DEBUG(dbgs() << r.toString() << "\n"); - repBytes += r.getLength(); +#if LLVM_VERSION_MAJOR > 3 + Replacements& replacements = Tool.getReplacements().begin()->second; +#else + Replacements& replacements = Tool.getReplacements(); +#endif + for (const auto &replacement : replacements) { + DEBUG(dbgs() << replacement.toString() << "\n"); + repBytes += replacement.getLength(); } std::ifstream src_file(dst, std::ios::binary | std::ios::ate); src_file.clear(); From d0feaee4f4c230e00daedcaea2c3f1ecd957cc3f Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 11:51:35 +0100 Subject: [PATCH 19/46] Avoid a double-free of HipifyPPCallbacks instance This bug was present all along, but something changed in the order of de-initialisation performed by llvm that makes it actually crash now. The constructor of HipifyPPCallbacks gives: ``` std::unique_ptr(this) ``` to the LLVM Preprocessor instance. The Preprocessor instance subsequently frees the HipifyPPCallbacks, which is then freed again when we leave the stack frame at line 4340. So: let's leak the HipifyPPCallbacks onto the heap, and leave the LLVM Preprocessor object responsible for tidying it up. [ROCm/clr commit: 893ee6d6ca12f353f6b16846c39411b587e552dc] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 8b16e75db0..0f73859893 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -4302,12 +4302,12 @@ int main(int argc, const char **argv) { replacementsToUse = &Tool.getReplacements(); #endif - HipifyPPCallbacks PPCallbacks(replacementsToUse, tmpFile); - Cuda2HipCallback Callback(replacementsToUse, &Finder, &PPCallbacks, tmpFile); + HipifyPPCallbacks* PPCallbacks = new HipifyPPCallbacks(replacementsToUse, tmpFile); + Cuda2HipCallback Callback(replacementsToUse, &Finder, PPCallbacks, tmpFile); addAllMatchers(Finder, &Callback); - auto action = newFrontendActionFactory(&Finder, &PPCallbacks); + auto action = newFrontendActionFactory(&Finder, PPCallbacks); Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("--cuda-host-only", ArgumentInsertPosition::BEGIN)); @@ -4365,13 +4365,13 @@ int main(int argc, const char **argv) { } std::remove(csv.c_str()); } - if (0 == printStats(csv, src, PPCallbacks, Callback, repBytes, bytes, lines, start)) { + if (0 == printStats(csv, src, *PPCallbacks, Callback, repBytes, bytes, lines, start)) { filesTranslated--; } start = std::chrono::steady_clock::now(); repBytesTotal += repBytes; bytesTotal += bytes; - changedLinesTotal += PPCallbacks.LOCs.size() + Callback.LOCs.size(); + changedLinesTotal += PPCallbacks->LOCs.size() + Callback.LOCs.size(); linesTotal += lines; } dst.clear(); From eee60317fa8c14ed8e5548a5ab3c5d5e7758095c Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 21 Oct 2017 07:47:32 -0500 Subject: [PATCH 20/46] Use 2X for bidir memory bandwidth calc [ROCm/clr commit: 21689e87106ee45f05b7d0c16cf7f88880d31632] --- .../hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index b3b0b3e4a6..2d67b04441 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -594,7 +594,7 @@ void RunBenchmark_Bidir(ResultDatabase &resultDB) " ms\n"; } - double speed = (double(sizeToBytes(thisSize)) / (1000*1000)) / t; + double speed = (double(sizeToBytes(2*thisSize)) / (1000*1000)) / t; char sizeStr[256]; sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); resultDB.AddResult(std::string("Bidir_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, "GB/sec", speed); From 941682c701afc4d3f14767dc5dacb0ba1006afec Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 23 Oct 2017 10:24:36 +0530 Subject: [PATCH 21/46] Make elfio headers private Change-Id: I3ba174bb46e84a75380207d93a0da6fe3703689e [ROCm/clr commit: 281c543bc303f472ac31bfaad9e0eb1e5346cff7] --- .../hipamd/{include/hip/hcc_detail => src}/elfio/elf_types.hpp | 0 .../clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio.hpp | 0 .../hipamd/{include/hip/hcc_detail => src}/elfio/elfio_dump.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_dynamic.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_header.hpp | 0 .../hipamd/{include/hip/hcc_detail => src}/elfio/elfio_note.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_relocation.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_section.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_segment.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_strings.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_symbols.hpp | 0 .../{include/hip/hcc_detail => src}/elfio/elfio_utils.hpp | 0 projects/clr/hipamd/src/hip_module.cpp | 2 +- 13 files changed, 1 insertion(+), 1 deletion(-) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elf_types.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_dump.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_dynamic.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_header.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_note.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_relocation.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_section.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_segment.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_strings.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_symbols.hpp (100%) rename projects/clr/hipamd/{include/hip/hcc_detail => src}/elfio/elfio_utils.hpp (100%) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elf_types.hpp b/projects/clr/hipamd/src/elfio/elf_types.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elf_types.hpp rename to projects/clr/hipamd/src/elfio/elf_types.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio.hpp b/projects/clr/hipamd/src/elfio/elfio.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio.hpp rename to projects/clr/hipamd/src/elfio/elfio.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dump.hpp b/projects/clr/hipamd/src/elfio/elfio_dump.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dump.hpp rename to projects/clr/hipamd/src/elfio/elfio_dump.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dynamic.hpp b/projects/clr/hipamd/src/elfio/elfio_dynamic.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_dynamic.hpp rename to projects/clr/hipamd/src/elfio/elfio_dynamic.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_header.hpp b/projects/clr/hipamd/src/elfio/elfio_header.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_header.hpp rename to projects/clr/hipamd/src/elfio/elfio_header.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_note.hpp b/projects/clr/hipamd/src/elfio/elfio_note.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_note.hpp rename to projects/clr/hipamd/src/elfio/elfio_note.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_relocation.hpp b/projects/clr/hipamd/src/elfio/elfio_relocation.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_relocation.hpp rename to projects/clr/hipamd/src/elfio/elfio_relocation.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_section.hpp b/projects/clr/hipamd/src/elfio/elfio_section.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_section.hpp rename to projects/clr/hipamd/src/elfio/elfio_section.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_segment.hpp b/projects/clr/hipamd/src/elfio/elfio_segment.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_segment.hpp rename to projects/clr/hipamd/src/elfio/elfio_segment.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_strings.hpp b/projects/clr/hipamd/src/elfio/elfio_strings.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_strings.hpp rename to projects/clr/hipamd/src/elfio/elfio_strings.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_symbols.hpp b/projects/clr/hipamd/src/elfio/elfio_symbols.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_symbols.hpp rename to projects/clr/hipamd/src/elfio/elfio_symbols.hpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_utils.hpp b/projects/clr/hipamd/src/elfio/elfio_utils.hpp similarity index 100% rename from projects/clr/hipamd/include/hip/hcc_detail/elfio/elfio_utils.hpp rename to projects/clr/hipamd/src/elfio/elfio_utils.hpp diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index b6bc9d63fd..dee69e7ba0 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -34,7 +34,7 @@ THE SOFTWARE. #include #include -#include "hip/hcc_detail/elfio/elfio.hpp" +#include "elfio/elfio.hpp" #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" #include "trace_helper.h" From 49e43a7fc1ece49027e0d5d13c8aadbd59b03f2c Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:00:58 +0100 Subject: [PATCH 22/46] Remove dependency on nonexistent `lit` target This just... isn't a target. Maybe you were trying to do a file dependency? But that's handled implicitly by the find_program call anyway, so all this is doing is generating a cmake warning [ROCm/clr commit: e9d259699c96d723e8cd13a07cedb69a2972692f] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 355b0cea7a..d309aa261c 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -84,7 +84,7 @@ if (HIPIFY_CLANG_TESTS) add_lit_testsuite(test-hipify "Running HIPify regression tests" ${CMAKE_SOURCE_DIR}/tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit + DEPENDS hipify-clang ) add_custom_target(test-hipify-clang) From baabd2755eac0bae425b92c6763228840746cf0a Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:16:26 +0100 Subject: [PATCH 23/46] Propagate the CUDA toolkit directory into the lit tests Allows the tests to actually run... :D [ROCm/clr commit: 9747578d093164d7e5325aa1ef79958879e17b35] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 5 ++++- projects/clr/hipamd/tests/hipify-clang/axpy.cu | 2 +- projects/clr/hipamd/tests/hipify-clang/lit.cfg | 2 ++ projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in | 1 + 4 files changed, 8 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index d309aa261c..18939c0210 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -70,9 +70,12 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DI install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) - # tests find_package(PythonInterp 2.7 REQUIRED EXACT) + # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the + # value of --cuda-path for the test runs. + find_package(CUDA REQUIRED) + set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) configure_file( diff --git a/projects/clr/hipamd/tests/hipify-clang/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/axpy.cu index 8472e60209..56a854fd0a 100644 --- a/projects/clr/hipamd/tests/hipify-clang/axpy.cu +++ b/projects/clr/hipamd/tests/hipify-clang/axpy.cu @@ -1,4 +1,4 @@ -// RUN: hipify "%s" -o=%t -- +// RUN: hipify "%s" -o=%t -- %cuda_args #include diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index c57b8ec524..57d9d876b5 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/tests/hipify-clang/lit.cfg @@ -46,3 +46,5 @@ if obj_root is not None: config.substitutions.append(("hipify", obj_root+"/hipify-clang")) +# Clang args for CUDA... +config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30" % config.cuda_root)) diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in b/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in index 4511316ac7..4e05710065 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in +++ b/projects/clr/hipamd/tests/hipify-clang/lit.site.cfg.in @@ -2,6 +2,7 @@ import sys config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@BINARY_DIR@" +config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" # Support substitution of the tools and libs dirs with user parameters. This is # used when we can't determine the tool dir at configuration time. From d92c43bd21f9c533d625e09f231b2f5e099df431 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:27:22 +0100 Subject: [PATCH 24/46] Look for FileCheck for running lit tests, too Use of grep in `lit` RUN lines is deprecated: https://llvm.org/docs/TestingGuide.html#writing-new-regression-tests Using grep leads to really unhelpful failure output (it literally just says "the test failed"). FileCheck is much more helpful, and distributed with LLVM on most distros anyway, so this extra dependency shouldn't prove problematic. [ROCm/clr commit: 3868036ea7c80f46e9f305775d1145b883b5d0bb] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 18939c0210..e7c46557f8 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -3,11 +3,12 @@ project(hipify-clang) option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) -# Disable the tests if `lit` is not installed. +# Disable the tests if `lit` or `FileCheck` is not installed. find_program(LIT_COMMAND lit) -if (NOT LIT_COMMAND) +find_program(FILECHECK_COMMAND FileCheck) +if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND) set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") - message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") + message(STATUS "hipify-clang's tests are not being built because `lit` or `FileCheck` could not be found.") endif() set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") From 0f6c153774b4890404445c32b14b0c90e56d0ba3 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:45:16 +0100 Subject: [PATCH 25/46] Migrate lit test to using FileCheck, so failures are readable It seems the test is already broken, but look how awesome the error message is now: /home/chris/HIP/tests/hipify-clang/axpy.cu:31:12: error: expected string not found in input // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); ^ :31:2: note: scanning from here // ^ :33:2: note: possible intended match here hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); ^ [ROCm/clr commit: 74fd64d5c132081b56301e07b1220a608d8d8ac0] --- .../clr/hipamd/tests/hipify-clang/axpy.cu | 22 +++++++++++-------- 1 file changed, 13 insertions(+), 9 deletions(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/axpy.cu index 56a854fd0a..e0afa02a26 100644 --- a/projects/clr/hipamd/tests/hipify-clang/axpy.cu +++ b/projects/clr/hipamd/tests/hipify-clang/axpy.cu @@ -1,9 +1,9 @@ -// RUN: hipify "%s" -o=%t -- %cuda_args +// RUN: hipify "%s" -o=%t -- %cuda_args && cat %t | sed -Ee 's|//.+|// |g' | FileCheck %s #include __global__ void axpy(float a, float* x, float* y) { - // RUN: sh -c "test `grep -c -F 'y[hipThreadIdx_x] = a * x[hipThreadIdx_x];' %t` -eq 2" + // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } @@ -17,21 +17,25 @@ int main(int argc, char* argv[]) { // Copy input data to device. float* device_x; float* device_y; - // RUN: sh -c "test `grep -c -F 'hipMalloc(&device_x, kDataLen * sizeof(float));' %t` -eq 2" + + // CHECK: hipMalloc(&device_x, kDataLen * sizeof(float)); cudaMalloc(&device_x, kDataLen * sizeof(float)); - // RUN: sh -c "test `grep -c -F 'hipMalloc(&device_y, kDataLen * sizeof(float));' %t` -eq 2" + + // CHECK: hipMalloc(&device_y, kDataLen * sizeof(float)); cudaMalloc(&device_y, kDataLen * sizeof(float)); - // RUN: sh -c "test `grep -c -F 'hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);' %t` -eq 2" + + // CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice); cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel. - // RUN: sh -c "test `grep -c -F 'hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);' %t` -eq 2" + // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); // Copy output data to host. - // RUN: sh -c "test `grep -c -F 'hipDeviceSynchronize();' %t` -eq 2" + // CHECK: hipDeviceSynchronize(); cudaDeviceSynchronize(); - // RUN: sh -c "test `grep -c -F 'hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost);' %t` -eq 2" + + // CHECK: hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost); cudaMemcpy(host_y, device_y, kDataLen * sizeof(float), cudaMemcpyDeviceToHost); // Print the results. @@ -39,7 +43,7 @@ int main(int argc, char* argv[]) { std::cout << "y[" << i << "] = " << host_y[i] << "\n"; } - // RUN: sh -c "test `grep -c -F 'hipDeviceReset();' %t` -eq 2" + // CHECK: hipDeviceReset(); cudaDeviceReset(); return 0; } From 4e80782cda54b1e1d95ef1ac960479ac1a8481f6 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:46:36 +0100 Subject: [PATCH 26/46] Adapt `lit` test for the hipLaunchKernelGGL changes from before... [ROCm/clr commit: 5912f465bdc9bed19e3e6dbf7872e73de7c60383] --- projects/clr/hipamd/tests/hipify-clang/axpy.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/axpy.cu index e0afa02a26..689c6e96ff 100644 --- a/projects/clr/hipamd/tests/hipify-clang/axpy.cu +++ b/projects/clr/hipamd/tests/hipify-clang/axpy.cu @@ -28,7 +28,7 @@ int main(int argc, char* argv[]) { cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel. - // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); // Copy output data to host. From 8b311424a6d616236c81b395625060c38041c1de Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 13:50:00 +0100 Subject: [PATCH 27/46] Introduce a test runner script to simplify invocation ... And to use a standard, highly amusing trick for making coloured output work. [ROCm/clr commit: c99dcbba8d5e31c64098bb707f892d193e4da322] --- .../clr/hipamd/hipify-clang/CMakeLists.txt | 6 ++-- .../clr/hipamd/tests/hipify-clang/axpy.cu | 2 +- .../clr/hipamd/tests/hipify-clang/lit.cfg | 1 + .../clr/hipamd/tests/hipify-clang/run_test.sh | 28 +++++++++++++++++++ 4 files changed, 34 insertions(+), 3 deletions(-) create mode 100755 projects/clr/hipamd/tests/hipify-clang/run_test.sh diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index e7c46557f8..2ec96e51a8 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -6,9 +6,10 @@ option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed # Disable the tests if `lit` or `FileCheck` is not installed. find_program(LIT_COMMAND lit) find_program(FILECHECK_COMMAND FileCheck) -if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND) +find_program(SOCAT_COMMAND socat) +if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND OR NOT SOCAT_COMMAND) set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") - message(STATUS "hipify-clang's tests are not being built because `lit` or `FileCheck` could not be found.") + message(STATUS "hipify-clang's tests are not being built because `lit`,`FileCheck` or `socat` could not be found.") endif() set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") @@ -88,6 +89,7 @@ if (HIPIFY_CLANG_TESTS) add_lit_testsuite(test-hipify "Running HIPify regression tests" ${CMAKE_SOURCE_DIR}/tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ARGS -v DEPENDS hipify-clang ) diff --git a/projects/clr/hipamd/tests/hipify-clang/axpy.cu b/projects/clr/hipamd/tests/hipify-clang/axpy.cu index 689c6e96ff..8c6b0e0d8d 100644 --- a/projects/clr/hipamd/tests/hipify-clang/axpy.cu +++ b/projects/clr/hipamd/tests/hipify-clang/axpy.cu @@ -1,4 +1,4 @@ -// RUN: hipify "%s" -o=%t -- %cuda_args && cat %t | sed -Ee 's|//.+|// |g' | FileCheck %s +// RUN: %run_test hipify "%s" "%t" %cuda_args #include diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index 57d9d876b5..bb6ac2b407 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/tests/hipify-clang/lit.cfg @@ -48,3 +48,4 @@ config.substitutions.append(("hipify", obj_root+"/hipify-clang")) # Clang args for CUDA... config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30" % config.cuda_root)) +config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh")) diff --git a/projects/clr/hipamd/tests/hipify-clang/run_test.sh b/projects/clr/hipamd/tests/hipify-clang/run_test.sh new file mode 100755 index 0000000000..46b2fc066b --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/run_test.sh @@ -0,0 +1,28 @@ +#!/usr/bin/env bash + +set -o errexit + +# Run a single LIT test file in a magical way that preserves colour output, to work around +# a known flaw in lit. + +# Capture lit substitutions +HIPIFY=$1 +IN_FILE=$2 +TMP_FILE=$3 +shift 3 + +# Remaining args are the ones to forward to clang proper. + +# Time for the classic insane little trick for making colour output work. +# A self-deleting shell-script that does the thing we want to do... +TMP_SCRIPT=$(mktemp) +cat << EOF > $TMP_SCRIPT +set -o errexit +set -o xtrace +rm $TMP_SCRIPT +$HIPIFY -o=$TMP_FILE $IN_FILE -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE +EOF +chmod a+x $TMP_SCRIPT + +# Run the script via socat, spawning a virtual terminal and propagating exit code, and hence failure. +socat -du EXEC:$TMP_SCRIPT,pty,stderr STDOUT From 8fd3b3b1cd2207cb9bd412cfcba1927aee16b6ab Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 13:58:34 +0100 Subject: [PATCH 28/46] Add square.cu to lit testsuite [ROCm/clr commit: ead79e5bf4e54f0f7b1fda3482c2d1b0a36d8022] --- .../clr/hipamd/tests/hipify-clang/square.cu | 114 ++++++++++++++++++ 1 file changed, 114 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/square.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/square.cu b/projects/clr/hipamd/tests/hipify-clang/square.cu new file mode 100644 index 0000000000..e0c72094a8 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/square.cu @@ -0,0 +1,114 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include +#include + +#define CHECK(cmd) \ +{\ + cudaError_t error = cmd;\ + if (error != cudaSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + + +/* + * Square each element in the array A and write to array C. + */ +template +__global__ void +vector_square(T *C_d, const T *A_d, size_t N) +{ + // CHECK: size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + // CHECK: size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i=offset; i>> (C_d, A_d, N); + + printf ("info: copy Device2Host\n"); + // CHECK: CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + CHECK ( cudaMemcpy(C_h, C_d, Nbytes, cudaMemcpyDeviceToHost)); + + printf ("info: check result\n"); + for (size_t i=0; i Date: Wed, 18 Oct 2017 14:05:46 +0100 Subject: [PATCH 29/46] Add cudaRegister.cu lit test [ROCm/clr commit: 2437b319398daa36d32b2e083a39e6c3cd30fe00] --- .../hipamd/tests/hipify-clang/cudaRegister.cu | 111 ++++++++++++++++++ 1 file changed, 111 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu b/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu new file mode 100644 index 0000000000..80d17f65b9 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/cudaRegister.cu @@ -0,0 +1,111 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include +#include +#include +#include + +#define LEN 1024 +#define SIZE LEN * sizeof(float) +#define ITER 1024*1024 + +// CHECK: if(status != hipSuccess) { +#define check(msg, status){ \ +if(status != cudaSuccess) { \ + printf("%s failed. \n", #msg); \ +} \ +} + +__global__ void Inc1(float *Ad, float *Bd){ + // CHECK: int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if(tx < 1 ){ + for(int i=0;i>>(Ad, Bd); + sleep(3); + A[0] = -(ITER*1.0f); + std::cout<<"Same cache line before completion: \t"<< A[0]<>>(Ad, Bd); + sleep(3); + A[0] = -(ITER*1.0f); + std::cout<<"Diff cache line before completion: \t"< Date: Wed, 18 Oct 2017 20:41:23 +0100 Subject: [PATCH 30/46] Add the CUDA samples include dir to the path for tests Means we get to easily steal CUDA examples for tests [ROCm/clr commit: 64d5f07050165e15325a3791540b10efad6cce43] --- projects/clr/hipamd/tests/hipify-clang/lit.cfg | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index bb6ac2b407..76b2ca08dc 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/tests/hipify-clang/lit.cfg @@ -47,5 +47,5 @@ if obj_root is not None: config.substitutions.append(("hipify", obj_root+"/hipify-clang")) # Clang args for CUDA... -config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30" % config.cuda_root)) +config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30 -isystem%s/samples/common/inc" % (config.cuda_root, config.cuda_root))) config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh")) From 26be50b0826abf8119fb3e74fe61481692e97e54 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 21:12:35 +0100 Subject: [PATCH 31/46] Add concurentKernels.cu to the testsuite [ROCm/clr commit: 2c65d0da37540e27d67cdd2b357f5d5e5a0aa22e] --- .../tests/hipify-clang/concurentKernels.cu | 240 ++++++++++++++++++ 1 file changed, 240 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu b/projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu new file mode 100644 index 0000000000..e369baaf3e --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/concurentKernels.cu @@ -0,0 +1,240 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +/* + * Copyright 1993-2015 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +// +// This sample demonstrates the use of streams for concurrent execution. It also illustrates how to +// introduce dependencies between CUDA streams with the new cudaStreamWaitEvent function introduced +// in CUDA 3.2. +// +// Devices of compute capability 1.x will run the kernels one after another +// Devices of compute capability 2.0 or higher can overlap the kernels +// +#include +#include +#include + +// This is a kernel that does no real work but runs at least for a specified number of clocks +__global__ void clock_block(clock_t *d_o, clock_t clock_count) +{ + unsigned int start_clock = (unsigned int) clock(); + + clock_t clock_offset = 0; + + while (clock_offset < clock_count) + { + unsigned int end_clock = (unsigned int) clock(); + + // The code below should work like + // this (thanks to modular arithmetics): + // + // clock_offset = (clock_t) (end_clock > start_clock ? + // end_clock - start_clock : + // end_clock + (0xffffffffu - start_clock)); + // + // Indeed, let m = 2^32 then + // end - start = end + m - start (mod m). + + clock_offset = (clock_t)(end_clock - start_clock); + } + + d_o[0] = clock_offset; +} + + +// Single warp reduction kernel +__global__ void sum(clock_t *d_clocks, int N) +{ + __shared__ clock_t s_clocks[32]; + + clock_t my_sum = 0; + + for (int i = threadIdx.x; i < N; i+= blockDim.x) + { + my_sum += d_clocks[i]; + } + + s_clocks[threadIdx.x] = my_sum; + syncthreads(); + + for (int i=16; i>0; i/=2) + { + if (threadIdx.x < i) + { + s_clocks[threadIdx.x] += s_clocks[threadIdx.x + i]; + } + + syncthreads(); + } + + d_clocks[0] = s_clocks[0]; +} + +int main(int argc, char **argv) +{ + int nkernels = 8; // number of concurrent kernels + int nstreams = nkernels + 1; // use one more stream than concurrent kernel + int nbytes = nkernels * sizeof(clock_t); // number of data bytes + float kernel_time = 10; // time the kernel should run in ms + float elapsed_time; // timing variables + int cuda_device = 0; + + printf("[%s] - Starting...\n", argv[0]); + + // get number of kernels if overridden on the command line + if (checkCmdLineFlag(argc, (const char **)argv, "nkernels")) + { + nkernels = getCmdLineArgumentInt(argc, (const char **)argv, "nkernels"); + nstreams = nkernels + 1; + } + + // use command-line specified CUDA device, otherwise use device with highest Gflops/s + cuda_device = findCudaDevice(argc, (const char **)argv); + + cudaDeviceProp deviceProp; + // CHECK: checkCudaErrors(hipGetDevice(&cuda_device)); + checkCudaErrors(cudaGetDevice(&cuda_device)); + + // CHECK: checkCudaErrors(hipGetDeviceProperties(&deviceProp, cuda_device)); + checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device)); + + if ((deviceProp.concurrentKernels == 0)) + { + printf("> GPU does not support concurrent kernel execution\n"); + printf(" CUDA kernel runs will be serialized\n"); + } + + printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n", + deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount); + + // allocate host memory + clock_t *a = 0; // pointer to the array data in host memory + // CHECK: checkCudaErrors(hipHostMalloc((void **)&a, nbytes)); + checkCudaErrors(cudaMallocHost((void **)&a, nbytes)); + + // allocate device memory + clock_t *d_a = 0; // pointers to data and init value in the device memory + // CHECK: checkCudaErrors(hipMalloc((void **)&d_a, nbytes)); + checkCudaErrors(cudaMalloc((void **)&d_a, nbytes)); + + // CHECK: hipStream_t *streams = (hipStream_t *) malloc(nstreams * sizeof(hipStream_t)); + // allocate and initialize an array of stream handles + cudaStream_t *streams = (cudaStream_t *) malloc(nstreams * sizeof(cudaStream_t)); + + for (int i = 0; i < nstreams; i++) + { + // CHECK: checkCudaErrors(hipStreamCreate(&(streams[i]))); + checkCudaErrors(cudaStreamCreate(&(streams[i]))); + } + + // create CUDA event handles + cudaEvent_t start_event, stop_event; + + // CHECK: checkCudaErrors(hipEventCreate(&start_event)); + // CHECK: checkCudaErrors(hipEventCreate(&stop_event)); + checkCudaErrors(cudaEventCreate(&start_event)); + checkCudaErrors(cudaEventCreate(&stop_event)); + + // the events are used for synchronization only and hence do not need to record timings + // this also makes events not introduce global sync points when recorded which is critical to get overlap + + // CHECK: hipEvent_t *kernelEvent; + // CHECK: kernelEvent = (hipEvent_t *) malloc(nkernels * sizeof(hipEvent_t)); + cudaEvent_t *kernelEvent; + kernelEvent = (cudaEvent_t *) malloc(nkernels * sizeof(cudaEvent_t)); + + for (int i = 0; i < nkernels; i++) + { + // CHECK: checkCudaErrors(hipEventCreateWithFlags(&(kernelEvent[i]), hipEventDisableTiming)); + checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming)); + } + + ////////////////////////////////////////////////////////////////////// + // time execution with nkernels streams + clock_t total_clocks = 0; +#if defined(__arm__) || defined(__aarch64__) + // the kernel takes more time than the channel reset time on arm archs, so to prevent hangs reduce time_clocks. + clock_t time_clocks = (clock_t)(kernel_time * (deviceProp.clockRate / 1000)); +#else + clock_t time_clocks = (clock_t)(kernel_time * deviceProp.clockRate); +#endif + + // CHECK: hipEventRecord(start_event, 0); + cudaEventRecord(start_event, 0); + + // queue nkernels in separate streams and record when they are done + for (int i=0; i>>(&d_a[i], time_clocks); + total_clocks += time_clocks; + + // CHECK: checkCudaErrors(hipEventRecord(kernelEvent[i], streams[i])); + checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i])); + + // make the last stream wait for the kernel event to be recorded + // CHECK: checkCudaErrors(hipStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); + checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); + } + + // queue a sum kernel and a copy back to host in the last stream. + // the commands in this stream get dispatched as soon as all the kernel events have been recorded + // CHECK: hipLaunchKernelGGL(sum, dim3(1), dim3(32), 0, streams[nstreams-1], d_a, nkernels); + // CHECK: checkCudaErrors(hipMemcpyAsync(a, d_a, sizeof(clock_t), hipMemcpyDeviceToHost, streams[nstreams-1])); + sum<<<1,32,0,streams[nstreams-1]>>>(d_a, nkernels); + checkCudaErrors(cudaMemcpyAsync(a, d_a, sizeof(clock_t), cudaMemcpyDeviceToHost, streams[nstreams-1])); + + // at this point the CPU has dispatched all work for the GPU and can continue processing other tasks in parallel + + // in this sample we just wait until the GPU is done + // CHECK: checkCudaErrors(hipEventRecord(stop_event, 0)); + // CHECK: checkCudaErrors(hipEventSynchronize(stop_event)); + // CHECK: checkCudaErrors(hipEventElapsedTime(&elapsed_time, start_event, stop_event)); + checkCudaErrors(cudaEventRecord(stop_event, 0)); + checkCudaErrors(cudaEventSynchronize(stop_event)); + checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start_event, stop_event)); + + printf("Expected time for serial execution of %d kernels = %.3fs\n", nkernels, nkernels * kernel_time/1000.0f); + printf("Expected time for concurrent execution of %d kernels = %.3fs\n", nkernels, kernel_time/1000.0f); + printf("Measured time for sample = %.3fs\n", elapsed_time/1000.0f); + + bool bTestResult = (a[0] > total_clocks); + + // release resources + for (int i = 0; i < nkernels; i++) + { + // CHECK: hipStreamDestroy(streams[i]); + // CHECK: hipEventDestroy(kernelEvent[i]); + cudaStreamDestroy(streams[i]); + cudaEventDestroy(kernelEvent[i]); + } + + free(streams); + free(kernelEvent); + + // CHECK: hipEventDestroy(start_event); + // CHECK: hipEventDestroy(stop_event); + // CHECK: hipHostFree(a); + // CHECK: hipFree(d_a); + cudaEventDestroy(start_event); + cudaEventDestroy(stop_event); + cudaFreeHost(a); + cudaFree(d_a); + + if (!bTestResult) + { + printf("Test failed!\n"); + exit(EXIT_FAILURE); + } + + printf("Test passed\n"); + exit(EXIT_SUCCESS); +} From dc6c43772ceeaefed92fd47ec71541d531603c99 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 23 Oct 2017 21:57:20 +0530 Subject: [PATCH 32/46] Use 2X for bidir p2p memory bandwidth calc [ROCm/clr commit: 4090c82936328643b8e45edd6fc464ea7d7bd2e0] --- .../hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 2d67b04441..4bd1bc0163 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -922,7 +922,7 @@ void RunBenchmark_P2P_Bidir(ResultDatabase &resultDB) { std::cerr << "size " << sizeToString(thisSize) << " took " << t << " ms\n"; } - double speed = (double(sizeToBytes(thisSize) * p_beatsperiteration) / (1000*1000)) / t; + double speed = (double(sizeToBytes(2*thisSize) * p_beatsperiteration) / (1000*1000)) / t; char sizeStr[256]; if (p_beatsperiteration>1) { sprintf(sizeStr, "%9sx%d", sizeToString(thisSize).c_str(), p_beatsperiteration); From 27df5ecc829ef7acaa16e56ddb77c5895e85e4d2 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 23 Oct 2017 21:16:13 +0300 Subject: [PATCH 33/46] [HIPIFY] cmake: fix standalone build [ROCm/clr commit: cdd849541f6a5f17c6db4c16b8d432b82bd1033e] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 355b0cea7a..2d9cbb1feb 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -75,15 +75,21 @@ if (HIPIFY_CLANG_TESTS) set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + if (CMAKE_CURRENT_SOURCE_DIR EQUAL CMAKE_SOURCE_DIR) + set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}) + else() + set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/..) + endif() + configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ${SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${BINARY_DIR}/tests/hipify-clang/lit.site.cfg @ONLY ) add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ${SOURCE_DIR}/tests/hipify-clang + PARAMS site_config=${BINARY_DIR}/tests/hipify-clang/lit.site.cfg DEPENDS hipify-clang lit ) From 482541bb8c2886810c7372b04c653296104a5dca Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Tue, 24 Oct 2017 01:45:23 +0100 Subject: [PATCH 34/46] Tweak some version numbers in clang version compatibility checks Apparently a couple of those APIs changed in clang 5, not 4. Drat. [ROCm/clr commit: aa288fcf1fab4b7426ef88e3730cfbe65d63c6cc] --- projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp index 0f73859893..1f0bcdba65 100644 --- a/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3031,7 +3031,7 @@ public: : Cuda2Hip(R, mainFileName) {} virtual bool handleBeginSource(CompilerInstance &CI -#if LLVM_VERSION_MAJOR < 4 +#if LLVM_VERSION_MAJOR <= 4 , StringRef Filename #endif ) override { @@ -3114,7 +3114,7 @@ public: if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics // remain unchanged. -#if LLVM_VERSION_MAJOR > 3 +#if LLVM_VERSION_MAJOR > 4 #define GET_NUM_ARGS() getNumParams() #else #define GET_NUM_ARGS() getNumArgs() From d9a550d09c2e84937b922d8e1b38edff8a2d89f2 Mon Sep 17 00:00:00 2001 From: emankov Date: Tue, 24 Oct 2017 10:51:11 +0300 Subject: [PATCH 35/46] [HIPIFY] cmake: simplify build [ROCm/clr commit: 559d9a68aab1a40c3ca60d22303b3d84b2491a0e] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 2d9cbb1feb..5f4f96fb8a 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -73,23 +73,15 @@ if (HIPIFY_CLANG_TESTS) # tests find_package(PythonInterp 2.7 REQUIRED EXACT) - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) - - if (CMAKE_CURRENT_SOURCE_DIR EQUAL CMAKE_SOURCE_DIR) - set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}) - else() - set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/..) - endif() - configure_file( - ${SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${BINARY_DIR}/tests/hipify-clang/lit.site.cfg + $${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg @ONLY ) add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang + PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg DEPENDS hipify-clang lit ) From 10cd9f9eb954b705aed16585de0fc4a4344dfb1a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 24 Oct 2017 14:16:05 +0300 Subject: [PATCH 36/46] [HIPIFY][fix] cmake: do not build hipify-clang if not asked + warn "hipify-clang will not be built" if HIPIFY_CLANG_LLVM_DIR is not specified. + fix typo in previous commit . [ROCm/clr commit: 7c51376bbf0238885d21f07e6adab40d10c3a515] --- projects/clr/hipamd/hipify-clang/CMakeLists.txt | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/CMakeLists.txt b/projects/clr/hipamd/hipify-clang/CMakeLists.txt index 5f4f96fb8a..b9d009ad92 100644 --- a/projects/clr/hipamd/hipify-clang/CMakeLists.txt +++ b/projects/clr/hipamd/hipify-clang/CMakeLists.txt @@ -1,6 +1,15 @@ cmake_minimum_required(VERSION 2.8.8) project(hipify-clang) +set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") + +if (HIPIFY_CLANG_LLVM_DIR) + find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) +else() + message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or higher using HIPIFY_CLANG_LLVM_DIR") + return() +endif() + option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) # Disable the tests if `lit` is not installed. @@ -10,10 +19,6 @@ if (NOT LIT_COMMAND) message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") endif() -set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") - -find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) - list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) include(AddLLVM) @@ -74,7 +79,7 @@ if (HIPIFY_CLANG_TESTS) find_package(PythonInterp 2.7 REQUIRED EXACT) configure_file( - $${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg @ONLY ) From 28c72504b636d58d141abfd195d49cf0bf2eb9d5 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 24 Oct 2017 18:12:25 +0530 Subject: [PATCH 37/46] Example showing globals use with module APIs [ROCm/clr commit: 10b1b5850511eb0c579cf451baa1e125f46e06d2] --- .../0_Intro/module_api_global/Makefile | 17 ++ .../0_Intro/module_api_global/runKernel.cpp | 167 ++++++++++++++++++ .../0_Intro/module_api_global/vcpy_kernel.cpp | 41 +++++ 3 files changed, 225 insertions(+) create mode 100644 projects/clr/hipamd/samples/0_Intro/module_api_global/Makefile create mode 100644 projects/clr/hipamd/samples/0_Intro/module_api_global/runKernel.cpp create mode 100644 projects/clr/hipamd/samples/0_Intro/module_api_global/vcpy_kernel.cpp diff --git a/projects/clr/hipamd/samples/0_Intro/module_api_global/Makefile b/projects/clr/hipamd/samples/0_Intro/module_api_global/Makefile new file mode 100644 index 0000000000..97605020b2 --- /dev/null +++ b/projects/clr/hipamd/samples/0_Intro/module_api_global/Makefile @@ -0,0 +1,17 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif +HIPCC=$(HIP_PATH)/bin/hipcc +HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) + +all: vcpy_kernel.code runKernel.hip.out + +runKernel.hip.out: runKernel.cpp + $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + +vcpy_kernel.code: vcpy_kernel.cpp + $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ + +clean: + rm -f *.code *.out diff --git a/projects/clr/hipamd/samples/0_Intro/module_api_global/runKernel.cpp b/projects/clr/hipamd/samples/0_Intro/module_api_global/runKernel.cpp new file mode 100644 index 0000000000..3f84720edf --- /dev/null +++ b/projects/clr/hipamd/samples/0_Intro/module_api_global/runKernel.cpp @@ -0,0 +1,167 @@ +/* +Copyright (c) 2017 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include +#include + +#define LEN 64 +#define SIZE LEN*sizeof(float) + +#define fileName "vcpy_kernel.code" +float myDeviceGlobal; +float myDeviceGlobalArray[16]; +#define HIP_CHECK(cmd) \ +{\ + hipError_t status = cmd;\ + if(status != hipSuccess) {std::cout<<"error: #"<