From 6cedd3ef4973f3785757413db89a7c5d0ee2b58b Mon Sep 17 00:00:00 2001 From: Mahmoud Date: Fri, 22 May 2020 20:29:21 -0400 Subject: moving ISA def files to new folder --- src/trace-driven/ISA_Def/kepler_opcode.h | 149 +++++ src/trace-driven/ISA_Def/pascal_opcode.h | 201 +++++++ src/trace-driven/ISA_Def/trace_opcode.h | 71 +++ src/trace-driven/ISA_Def/turing_opcode.h | 24 + src/trace-driven/ISA_Def/volta_opcode.h | 175 ++++++ src/trace-driven/gpgpusim_trace_driven_main.cc | 2 +- src/trace-driven/kepler_opcode.h | 150 ----- src/trace-driven/pascal_opcode.h | 202 ------- src/trace-driven/trace_driven.cc | 10 +- src/trace-driven/trace_driven.h | 2 +- src/trace-driven/trace_opcode.h | 71 --- .../traces-generator/NVbit_tool/Makefile | 23 - .../traces-generator/NVbit_tool/traceall.cu | 646 --------------------- src/trace-driven/traces-generator/README | 1 - .../traces-post-processing/Makefile | 11 - .../post-traces-processing.cpp | 195 ------- src/trace-driven/turing_opcode.h | 25 - src/trace-driven/volta_opcode.h | 176 ------ 18 files changed, 627 insertions(+), 1507 deletions(-) create mode 100644 src/trace-driven/ISA_Def/kepler_opcode.h create mode 100644 src/trace-driven/ISA_Def/pascal_opcode.h create mode 100644 src/trace-driven/ISA_Def/trace_opcode.h create mode 100644 src/trace-driven/ISA_Def/turing_opcode.h create mode 100644 src/trace-driven/ISA_Def/volta_opcode.h delete mode 100644 src/trace-driven/kepler_opcode.h delete mode 100644 src/trace-driven/pascal_opcode.h delete mode 100644 src/trace-driven/trace_opcode.h delete mode 100644 src/trace-driven/traces-generator/NVbit_tool/Makefile delete mode 100644 src/trace-driven/traces-generator/NVbit_tool/traceall.cu delete mode 100644 src/trace-driven/traces-generator/README delete mode 100755 src/trace-driven/traces-generator/traces-post-processing/Makefile delete mode 100644 src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp delete mode 100644 src/trace-driven/turing_opcode.h delete mode 100644 src/trace-driven/volta_opcode.h diff --git a/src/trace-driven/ISA_Def/kepler_opcode.h b/src/trace-driven/ISA_Def/kepler_opcode.h new file mode 100644 index 0000000..675ea6c --- /dev/null +++ b/src/trace-driven/ISA_Def/kepler_opcode.h @@ -0,0 +1,149 @@ +//developed by Mahmoud Khairy, Purdue Univ +//abdallm@purdue.edu + +#ifndef KEPLER_OPCODE_H +#define KEPLER_OPCODE_H + +#include "trace_opcode.h" +#include +#include + +#define KEPLER_BINART_VERSION 35 +#define KEPLER_SHARED_MEMORY_VIRTIAL_ADDRESS_START 0x00007f2c60000000 + +//TO DO: moving this to a yml or def files + +///Kepler ISA +//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html +static const std::unordered_map Kepler_OpcodeMap = { + //Floating Point 32 Instructions + {"FFMA", OpcodeChar(OP_FFMA, SP_OP)}, + {"FFMA32I", OpcodeChar(OP_FFMA32I, SP_OP)}, + {"FADD", OpcodeChar(OP_FADD, SP_OP)}, + {"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)}, + {"FCMP", OpcodeChar(OP_FCMP, SP_OP)}, + {"FMUL", OpcodeChar(OP_FMUL, SP_OP)}, + {"FMUL32I", OpcodeChar(OP_FMUL32I, SP_OP)}, + {"FMNMX", OpcodeChar(OP_FMNMX, SP_OP)}, + {"FSWZ", OpcodeChar(OP_FSWZ, SP_OP)}, + {"FSET", OpcodeChar(OP_FSET, SP_OP)}, + {"FSETP", OpcodeChar(OP_FSETP, SP_OP)}, + {"FCHK", OpcodeChar(OP_FCHK, SP_OP)}, + {"RRO", OpcodeChar(OP_RRO, SP_OP)}, + //SFU + {"MUFU", OpcodeChar(OP_MUFU, SFU_OP)}, + + + //Double Point Instructions + {"DFMA", OpcodeChar(OP_DFMA, DP_OP)}, + {"DADD", OpcodeChar(OP_DADD, DP_OP)}, + {"DMUL", OpcodeChar(OP_DMUL, DP_OP)}, + {"DMNMX", OpcodeChar(OP_DMNMX, DP_OP)}, + {"DSET", OpcodeChar(OP_DSET, DP_OP)}, + {"DSETP", OpcodeChar(OP_DSETP, DP_OP)}, + + //Integer Instructions + {"IMAD", OpcodeChar(OP_IMAD, INTP_OP)}, + {"IMADSP", OpcodeChar(OP_IMADSP, INTP_OP)}, + {"IMUL", OpcodeChar(OP_IMUL, INTP_OP)}, + {"IMUL32I", OpcodeChar(OP_IMUL32I, INTP_OP)}, + {"IADD", OpcodeChar(OP_IADD, INTP_OP)}, + {"IADD32I", OpcodeChar(OP_IADD32I, INTP_OP)}, + {"ISUB", OpcodeChar(OP_ISUB, INTP_OP)}, + {"ISCADD", OpcodeChar(OP_ISCADD, INTP_OP)}, + {"ISCADD32I", OpcodeChar(OP_ISCADD32I, INTP_OP)}, + {"ISAD", OpcodeChar(OP_ISAD, INTP_OP)}, + {"IMNMX", OpcodeChar(OP_IMNMX, INTP_OP)}, + {"BFE", OpcodeChar(OP_BFE, INTP_OP)}, + {"BFI", OpcodeChar(OP_BFI, INTP_OP)}, + {"SHR", OpcodeChar(OP_SHR, INTP_OP)}, + {"SHL", OpcodeChar(OP_SHL, INTP_OP)}, + {"SHF", OpcodeChar(OP_SHF, INTP_OP)}, + {"LOP", OpcodeChar(OP_LOP, INTP_OP)}, + {"LOP32I", OpcodeChar(OP_LOP32I, INTP_OP)}, + {"FLO", OpcodeChar(OP_FLO, INTP_OP)}, + {"ISET", OpcodeChar(OP_ISET, INTP_OP)}, + {"ISETP", OpcodeChar(OP_ISETP, INTP_OP)}, + {"ICMP", OpcodeChar(OP_ICMP, INTP_OP)}, + {"POPC", OpcodeChar(OP_POPC, INTP_OP)}, + + //Conversion Instructions + {"F2F", OpcodeChar(OP_F2F, ALU_OP)}, + {"F2I", OpcodeChar(OP_F2I, ALU_OP)}, + {"I2F", OpcodeChar(OP_I2F, ALU_OP)}, + {"I2I", OpcodeChar(OP_I2I, ALU_OP)}, + + //Movement Instructions + {"MOV", OpcodeChar(OP_MOV, ALU_OP)}, + {"MOV32I", OpcodeChar(OP_MOV32I, ALU_OP)}, + {"SEL", OpcodeChar(OP_SEL, ALU_OP)}, + {"PRMT", OpcodeChar(OP_PRMT, ALU_OP)}, + {"SHFL", OpcodeChar(OP_SHFL, ALU_OP)}, + + //Predicate Instructions + {"P2R", OpcodeChar(OP_P2R, ALU_OP)}, + {"R2P", OpcodeChar(OP_R2P, ALU_OP)}, + {"CSET", OpcodeChar(OP_CSET, ALU_OP)}, + {"CSETP", OpcodeChar(OP_CSETP, ALU_OP)}, + {"PSET", OpcodeChar(OP_PSET, ALU_OP)}, + {"PSETP", OpcodeChar(OP_PSETP, ALU_OP)}, + + //Texture Instructions + //For now, we ignore texture loads, consider it as ALU_OP + {"TEX", OpcodeChar(OP_TEX, ALU_OP)}, + {"TLD", OpcodeChar(OP_TLD, ALU_OP)}, + {"TLD4", OpcodeChar(OP_TLD4, ALU_OP)}, + {"TXQ", OpcodeChar(OP_TXQ, ALU_OP)}, + + //Load/Store Instructions + //For now, we ignore constant loads, consider it as ALU_OP, TO DO + {"LDC", OpcodeChar(OP_LDC, ALU_OP)}, + //in Kepler, LD is load global so set it to LDG + {"LD", OpcodeChar(OP_LDG, LOAD_OP)}, + {"LDG", OpcodeChar(OP_LDG, LOAD_OP)}, + {"LDL", OpcodeChar(OP_LDL, LOAD_OP)}, + {"LDS", OpcodeChar(OP_LDS, LOAD_OP)}, + {"LDSLK", OpcodeChar(OP_LDSLK, LOAD_OP)}, + {"ST", OpcodeChar(OP_STG, STORE_OP)}, + {"STL", OpcodeChar(OP_STL, STORE_OP)}, + {"STS", OpcodeChar(OP_STS, STORE_OP)}, + {"STSCUL", OpcodeChar(OP_STSCUL, STORE_OP)}, + {"ATOM", OpcodeChar(OP_ATOM, STORE_OP)}, + {"RED", OpcodeChar(OP_RED, STORE_OP)}, + {"CCTL", OpcodeChar(OP_CCTL, ALU_OP)}, + {"CCTLL", OpcodeChar(OP_CCTLL, ALU_OP)}, + {"MEMBAR", OpcodeChar(OP_MEMBAR, MEMORY_BARRIER_OP)}, + + //surface memory instructions + {"SUCLAMP", OpcodeChar(OP_SUCLAMP, LOAD_OP)}, + {"SUBFM", OpcodeChar(OP_SUBFM, LOAD_OP)}, + {"SUEAU", OpcodeChar(OP_SUEAU, LOAD_OP)}, + {"SULDGA", OpcodeChar(OP_SULDGA, LOAD_OP)}, + {"SUSTGA", OpcodeChar(OP_SUSTGA, STORE_OP)}, + + //Control Instructions + {"BRA", OpcodeChar(OP_BRA, BRANCH_OP)}, + {"BRX", OpcodeChar(OP_BRX, BRANCH_OP)}, + {"JMP", OpcodeChar(OP_JMP, BRANCH_OP)}, + {"JMX", OpcodeChar(OP_JMX, BRANCH_OP)}, + {"CAL", OpcodeChar(OP_CAL, CALL_OPS)}, + {"JCAL", OpcodeChar(OP_JCAL, CALL_OPS)}, + {"RET", OpcodeChar(OP_RET, RET_OPS)}, + {"BRK", OpcodeChar(OP_BRK, RET_OPS)}, + {"CONT", OpcodeChar(OP_CONT, RET_OPS)}, + {"SSY", OpcodeChar(OP_SSY, RET_OPS)}, + {"PBK", OpcodeChar(OP_PBK, RET_OPS)}, + {"PCNT", OpcodeChar(OP_PCNT, RET_OPS)}, + {"PRET", OpcodeChar(OP_PRET, RET_OPS)}, + {"BPT", OpcodeChar(OP_BPT, BRANCH_OP)}, + {"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)}, + + //Miscellaneous Instructions + {"NOP", OpcodeChar(OP_NOP, ALU_OP)}, + {"S2R", OpcodeChar(OP_S2R, ALU_OP)}, + {"B2R", OpcodeChar(OP_B2R, ALU_OP)}, + {"BAR", OpcodeChar(OP_BAR, BARRIER_OP)}, + {"VOTE", OpcodeChar(OP_VOTE, ALU_OP)}, +}; + +#endif diff --git a/src/trace-driven/ISA_Def/pascal_opcode.h b/src/trace-driven/ISA_Def/pascal_opcode.h new file mode 100644 index 0000000..66a0841 --- /dev/null +++ b/src/trace-driven/ISA_Def/pascal_opcode.h @@ -0,0 +1,201 @@ +//developed by Mahmoud Khairy, Purdue Univ +//abdallm@purdue.edu + +#ifndef PASCAL_OPCODE_H +#define PASCAL_OPCODE_H + +#include "trace_opcode.h" +#include +#include + +#define PASCAL_TITANX_BINART_VERSION 61 +#define PASCAL_P100_BINART_VERSION 60 + +#define PASCAL_SHARED_MEMORY_VIRTIAL_ADDRESS_START 0x00007f2c60000000 + +//TO DO: moving this to a yml or def files + +///Pascal SM_61 ISA +//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html +static const std::unordered_map Pascal_OpcodeMap = { + //Floating Point 32 Instructions + {"FADD", OpcodeChar(OP_FADD, SP_OP)}, + {"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)}, + {"FCHK", OpcodeChar(OP_FCHK, SP_OP)}, + {"FFMA32I", OpcodeChar(OP_FFMA32I, SP_OP)}, + {"FFMA", OpcodeChar(OP_FFMA, SP_OP)}, + {"FMNMX", OpcodeChar(OP_FMNMX, SP_OP)}, + {"FMUL", OpcodeChar(OP_FMUL, SP_OP)}, + {"FMUL32I", OpcodeChar(OP_FMUL32I, SP_OP)}, + {"FSEL", OpcodeChar(OP_FSEL, SP_OP)}, + {"FSET", OpcodeChar(OP_FSET, SP_OP)}, + {"FSETP", OpcodeChar(OP_FSETP, SP_OP)}, + {"FSWZADD", OpcodeChar(OP_FSWZADD, SP_OP)}, + {"RRO", OpcodeChar(OP_RRO, SP_OP)}, + + //SFU + {"MUFU", OpcodeChar(OP_MUFU, SFU_OP)}, + + //Floating Point 16 Instructions + {"HADD2", OpcodeChar(OP_HADD2, SP_OP)}, + {"HFMA2", OpcodeChar(OP_HFMA2, SP_OP)}, + {"HMUL2", OpcodeChar(OP_HMUL2, SP_OP)}, + {"HSET2", OpcodeChar(OP_HSET2, SP_OP)}, + {"HSETP2", OpcodeChar(OP_HSETP2, SP_OP)}, + + //Double Point Instructions + {"DADD", OpcodeChar(OP_DADD, DP_OP)}, + {"DFMA", OpcodeChar(OP_DFMA, DP_OP)}, + {"DMUL", OpcodeChar(OP_DMUL, DP_OP)}, + {"DSETP", OpcodeChar(OP_DSETP, DP_OP)}, + {"DMNMX", OpcodeChar(OP_DMNMX, DP_OP)}, + {"DSET", OpcodeChar(OP_DSET, DP_OP)}, + + //Integer Instructions + {"BMSK", OpcodeChar(OP_BMSK, INTP_OP)}, + {"BREV", OpcodeChar(OP_BREV, INTP_OP)}, + {"FLO", OpcodeChar(OP_FLO, INTP_OP)}, + {"IABS", OpcodeChar(OP_IABS, INTP_OP)}, + {"IADD", OpcodeChar(OP_IADD, INTP_OP)}, + {"IADD3", OpcodeChar(OP_IADD3, INTP_OP)}, + {"IADD32I", OpcodeChar(OP_IADD32I, INTP_OP)}, + {"IDP", OpcodeChar(OP_IDP, INTP_OP)}, + {"IDP4A", OpcodeChar(OP_IDP4A, INTP_OP)}, + {"IMAD", OpcodeChar(OP_IMAD, INTP_OP)}, + {"IMMA", OpcodeChar(OP_IMMA, INTP_OP)}, + {"IMNMX", OpcodeChar(OP_IMNMX, INTP_OP)}, + {"IMUL", OpcodeChar(OP_IMUL, INTP_OP)}, + {"IMUL32I", OpcodeChar(OP_IMUL32I, INTP_OP)}, + {"ISCADD", OpcodeChar(OP_ISCADD, INTP_OP)}, + {"ISCADD32I", OpcodeChar(OP_ISCADD32I, INTP_OP)}, + {"ISETP", OpcodeChar(OP_ISETP, INTP_OP)}, + {"ISET", OpcodeChar(OP_ISET, INTP_OP)}, + {"LEA", OpcodeChar(OP_LEA, INTP_OP)}, + {"LOP", OpcodeChar(OP_LOP, INTP_OP)}, + {"LOP3", OpcodeChar(OP_LOP3, INTP_OP)}, + {"LOP32I", OpcodeChar(OP_LOP32I, INTP_OP)}, + {"POPC", OpcodeChar(OP_POPC, INTP_OP)}, + {"SHF", OpcodeChar(OP_SHF, INTP_OP)}, + {"SHR", OpcodeChar(OP_SHR, INTP_OP)}, + {"VABSDIFF", OpcodeChar(OP_VABSDIFF, INTP_OP)}, + {"VABSDIFF4", OpcodeChar(OP_VABSDIFF4, INTP_OP)}, + {"BFE", OpcodeChar(OP_BFE, INTP_OP)}, + {"BFI", OpcodeChar(OP_BFI, INTP_OP)}, + {"ICMP", OpcodeChar(OP_ICMP, INTP_OP)}, + {"IMADSP", OpcodeChar(OP_IMADSP, INTP_OP)}, + {"SHL", OpcodeChar(OP_SHL, INTP_OP)}, + {"XMAD", OpcodeChar(OP_XMAD, INTP_OP)}, + {"VMNMX", OpcodeChar(OP_VMNMX, INTP_OP)}, + + + //Conversion Instructions + {"F2F", OpcodeChar(OP_F2F, ALU_OP)}, + {"F2I", OpcodeChar(OP_F2I, ALU_OP)}, + {"I2F", OpcodeChar(OP_I2F, ALU_OP)}, + {"I2I", OpcodeChar(OP_I2I, ALU_OP)}, + {"I2IP", OpcodeChar(OP_I2IP, ALU_OP)}, + {"FRND", OpcodeChar(OP_FRND, ALU_OP)}, + + //Movement Instructions + {"MOV", OpcodeChar(OP_MOV, ALU_OP)}, + {"MOV32I", OpcodeChar(OP_MOV32I, ALU_OP)}, + {"PRMT", OpcodeChar(OP_PRMT, ALU_OP)}, + {"SEL", OpcodeChar(OP_SEL, ALU_OP)}, + {"SGXT", OpcodeChar(OP_SGXT, ALU_OP)}, + {"SHFL", OpcodeChar(OP_SHFL, ALU_OP)}, + + //Predicate Instructions + {"PLOP3", OpcodeChar(OP_PLOP3, ALU_OP)}, + {"PSETP", OpcodeChar(OP_PSETP, ALU_OP)}, + {"P2R", OpcodeChar(OP_P2R, ALU_OP)}, + {"R2P", OpcodeChar(OP_R2P, ALU_OP)}, + {"CSET", OpcodeChar(OP_CSET, ALU_OP)}, + {"CSETP", OpcodeChar(OP_CSETP, ALU_OP)}, + {"PSET", OpcodeChar(OP_PSET, ALU_OP)}, + + + //Load/Store Instructions + {"LD", OpcodeChar(OP_LD, LOAD_OP)}, + //For now, we ignore constant loads, consider it as ALU_OP, TO DO + {"LDC", OpcodeChar(OP_LDC, ALU_OP)}, + {"LDG", OpcodeChar(OP_LDG, LOAD_OP)}, + {"LDL", OpcodeChar(OP_LDL, LOAD_OP)}, + {"LDS", OpcodeChar(OP_LDS, LOAD_OP)}, + {"ST", OpcodeChar(OP_ST, STORE_OP)}, + {"STG", OpcodeChar(OP_STG, STORE_OP)}, + {"STL", OpcodeChar(OP_STL, STORE_OP)}, + {"STS", OpcodeChar(OP_STS, STORE_OP)}, + {"MATCH", OpcodeChar(OP_MATCH, ALU_OP)}, + {"QSPC", OpcodeChar(OP_QSPC, ALU_OP)}, + {"ATOM", OpcodeChar(OP_ATOM, STORE_OP)}, + {"ATOMS", OpcodeChar(OP_ATOMS, STORE_OP)}, + {"ATOMG", OpcodeChar(OP_ATOMG, STORE_OP)}, + {"RED", OpcodeChar(OP_RED, STORE_OP)}, + {"CCTL", OpcodeChar(OP_CCTL, ALU_OP)}, + {"CCTLL", OpcodeChar(OP_CCTLL, ALU_OP)}, + {"ERRBAR", OpcodeChar(OP_ERRBAR, ALU_OP)}, + {"MEMBAR", OpcodeChar(OP_MEMBAR, MEMORY_BARRIER_OP)}, + {"CCTLT", OpcodeChar(OP_CCTLT, ALU_OP)}, + + //Texture Instructions + //For now, we ignore texture loads, consider it as ALU_OP + {"TEX", OpcodeChar(OP_TEX, ALU_OP)}, + {"TLD", OpcodeChar(OP_TLD, ALU_OP)}, + {"TLD4", OpcodeChar(OP_TLD4, ALU_OP)}, + {"TMML", OpcodeChar(OP_TMML, ALU_OP)}, + {"TXD", OpcodeChar(OP_TXD, ALU_OP)}, + {"TXQ", OpcodeChar(OP_TXQ, ALU_OP)}, + {"TEXS", OpcodeChar(OP_TEXS, ALU_OP)}, + {"TLD4S", OpcodeChar(OP_TLD4S, ALU_OP)}, + {"TLDS", OpcodeChar(OP_TLDS, ALU_OP)}, + + //Control Instructions + {"BMOV", OpcodeChar(OP_BMOV, BRANCH_OP)}, + {"BPT", OpcodeChar(OP_BPT, BRANCH_OP)}, + {"BRA", OpcodeChar(OP_BRA, BRANCH_OP)}, + {"BREAK", OpcodeChar(OP_BREAK, BRANCH_OP)}, + {"BRX", OpcodeChar(OP_BRX, BRANCH_OP)}, + {"BSSY", OpcodeChar(OP_BSSY, BRANCH_OP)}, + {"BSYNC", OpcodeChar(OP_BSYNC, BRANCH_OP)}, + {"CALL", OpcodeChar(OP_CALL, CALL_OPS)}, + {"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)}, + {"JMP", OpcodeChar(OP_JMP, BRANCH_OP)}, + {"SSY", OpcodeChar(OP_SSY, BRANCH_OP)}, + {"SYNC", OpcodeChar(OP_SYNC, BRANCH_OP)}, + {"JMX", OpcodeChar(OP_JMX, BRANCH_OP)}, + {"KILL", OpcodeChar(OP_KILL, BRANCH_OP)}, + {"NANOSLEEP", OpcodeChar(OP_NANOSLEEP, BRANCH_OP)}, + {"RET", OpcodeChar(OP_RET, RET_OPS)}, + {"RPCMOV", OpcodeChar(OP_RPCMOV, BRANCH_OP)}, + {"RTT", OpcodeChar(OP_RTT, RET_OPS)}, + {"WARPSYNC", OpcodeChar(OP_WARPSYNC, BRANCH_OP)}, + {"YIELD", OpcodeChar(OP_YIELD, BRANCH_OP)}, + {"CAL", OpcodeChar(OP_CAL, CALL_OPS)}, + {"JCAL", OpcodeChar(OP_JCAL, CALL_OPS)}, + {"PRET", OpcodeChar(OP_PRET, CALL_OPS)}, + {"BRK", OpcodeChar(OP_BRK, CALL_OPS)}, + {"PBK", OpcodeChar(OP_PBK, CALL_OPS)}, + {"CONT", OpcodeChar(OP_CONT, CALL_OPS)}, + {"PCNT", OpcodeChar(OP_PCNT, CALL_OPS)}, + {"PEXIT", OpcodeChar(OP_PEXIT, CALL_OPS)}, + + //Miscellaneous Instructions + {"B2R", OpcodeChar(OP_B2R, ALU_OP)}, + {"BAR", OpcodeChar(OP_BAR, BARRIER_OP)}, + {"CS2R", OpcodeChar(OP_CS2R, ALU_OP)}, + {"CSMTEST", OpcodeChar(OP_CSMTEST, ALU_OP)}, + {"DEPBAR", OpcodeChar(OP_DEPBAR, ALU_OP)}, + {"GETLMEMBASE", OpcodeChar(OP_GETLMEMBASE, ALU_OP)}, + {"LEPC", OpcodeChar(OP_LEPC ,ALU_OP)}, + {"NOP", OpcodeChar(OP_NOP ,ALU_OP)}, + {"PMTRIG", OpcodeChar(OP_PMTRIG, ALU_OP)}, + {"R2B", OpcodeChar(OP_R2B, ALU_OP)}, + {"S2R", OpcodeChar(OP_S2R, ALU_OP)}, + {"SETCTAID", OpcodeChar(OP_SETCTAID, ALU_OP)}, + {"SETLMEMBASE", OpcodeChar(OP_SETLMEMBASE, ALU_OP)}, + {"VOTE", OpcodeChar(OP_VOTE, ALU_OP)}, + {"VOTE_VTG", OpcodeChar(OP_VOTE_VTG, ALU_OP)}, + +}; + +#endif diff --git a/src/trace-driven/ISA_Def/trace_opcode.h b/src/trace-driven/ISA_Def/trace_opcode.h new file mode 100644 index 0000000..ed147fc --- /dev/null +++ b/src/trace-driven/ISA_Def/trace_opcode.h @@ -0,0 +1,71 @@ +//developed by Mahmoud Khairy, Purdue Univ +//abdallm@purdue.edu + +#ifndef TRACE_OPCODE_H +#define TRACE_OPCODE_H + +#include "../../abstract_hardware_model.h" +#include +#include + + +enum TraceInstrOpcode { + //volta (common insts for others cards as well) + OP_FADD = 1, OP_FADD32I, OP_FCHK, OP_FFMA32I, OP_FFMA, OP_FMNMX, OP_FMUL, OP_FMUL32I, OP_FSEL, OP_FSET, OP_FSETP, + OP_FSWZADD, OP_MUFU, OP_HADD2, OP_HADD2_32I, OP_HFMA2, OP_HFMA2_32I, OP_HMUL2, OP_HMUL2_32I, OP_HSET2, OP_HSETP2, + OP_HMMA, OP_DADD, OP_DFMA, OP_DMUL, OP_DSETP, + OP_BMSK, OP_BREV, OP_FLO, OP_IABS, OP_IADD, OP_IADD3, OP_IADD32I, OP_IDP, OP_IDP4A, OP_IMAD, OP_IMMA, OP_IMNMX, + OP_IMUL, OP_IMUL32I, OP_ISCADD, OP_ISCADD32I, OP_ISETP, OP_LEA, OP_LOP, OP_LOP3, OP_LOP32I, OP_POPC, OP_SHF, OP_SHR, + OP_VABSDIFF, OP_VABSDIFF4, + OP_F2F, OP_F2I, OP_I2F, OP_I2I, OP_I2IP, OP_FRND, OP_MOV, OP_MOV32I, OP_PRMT, OP_SEL, OP_SGXT, OP_SHFL, OP_PLOP3, + OP_PSETP, OP_P2R, OP_R2P, OP_LD, OP_LDC, OP_LDG, OP_LDL, OP_LDS, OP_ST, OP_STG, OP_STL, OP_STS, OP_MATCH, OP_QSPC, + OP_ATOM, OP_ATOMS, OP_ATOMG, OP_RED, OP_CCTL, OP_CCTLL, OP_ERRBAR, OP_MEMBAR, OP_CCTLT, + OP_TEX, OP_TLD, OP_TLD4, + OP_TMML, OP_TXD, OP_TXQ, OP_BMOV, OP_BPT, OP_BRA, OP_BREAK, OP_BRX, OP_BSSY, OP_BSYNC, OP_CALL, OP_EXIT, OP_JMP, OP_JMX, + OP_KILL, OP_NANOSLEEP, OP_RET, OP_RPCMOV, OP_RTT, OP_WARPSYNC, OP_YIELD, OP_B2R, OP_BAR, OP_CS2R, OP_CSMTEST, OP_DEPBAR, + OP_GETLMEMBASE, OP_LEPC, OP_NOP, OP_PMTRIG, OP_R2B, OP_S2R, OP_SETCTAID, OP_SETLMEMBASE, OP_VOTE, OP_VOTE_VTG, + //unique insts for pascal + OP_RRO, OP_DMNMX, OP_DSET, OP_BFE, OP_BFI, OP_ICMP, OP_IMADSP, OP_SHL, OP_XMAD, OP_CSET, OP_CSETP, + OP_TEXS, OP_TLD4S, OP_TLDS, OP_CAL, OP_JCAL, OP_PRET, OP_BRK, OP_PBK, OP_CONT, OP_PCNT, OP_PEXIT, OP_SSY, OP_SYNC, OP_PSET + , OP_VMNMX, OP_ISET, + //unique insts for kepler + OP_FCMP, OP_FSWZ, OP_ISAD, OP_LDSLK, OP_STSCUL, OP_SUCLAMP, OP_SUBFM, OP_SUEAU, OP_SULDGA, OP_SUSTGA, OP_ISUB, + SASS_NUM_OPCODES /* The total number of opcodes. */ +}; +typedef enum TraceInstrOpcode sass_op_type; + +/* +enum uarch_op_t { + NO_OP=-1, + ALU_OP=1, + SFU_OP, + TENSOR_CORE_OP, + DP_OP, + SP_OP, + INTP_OP, + ALU_SFU_OP, + LOAD_OP, + TENSOR_CORE_LOAD_OP, + TENSOR_CORE_STORE_OP, + STORE_OP, + BRANCH_OP, + BARRIER_OP, + MEMORY_BARRIER_OP, + CALL_OPS, + RET_OPS +}; +typedef enum uarch_op_t op_type; + */ + +struct OpcodeChar +{ + OpcodeChar(unsigned m_opcode, unsigned m_opcode_category) { + opcode = m_opcode; + opcode_category = m_opcode_category; + } + unsigned opcode; + unsigned opcode_category; +}; + + +#endif diff --git a/src/trace-driven/ISA_Def/turing_opcode.h b/src/trace-driven/ISA_Def/turing_opcode.h new file mode 100644 index 0000000..4df44d9 --- /dev/null +++ b/src/trace-driven/ISA_Def/turing_opcode.h @@ -0,0 +1,24 @@ +//developed by Mahmoud Khairy, Purdue Univ +//abdallm@purdue.edu + +#ifndef TURING_OPCODE_H +#define TURING_OPCODE_H + +#include "trace_opcode.h" +#include +#include + +//TO DO: moving this to a yml or def files + + +#define TURING_BINART_VERSION 72 + +///Tuing SM_72 ISA +//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html +static const std::unordered_map Turing_OpcodeMap = { + +//TO fill + +}; + +#endif diff --git a/src/trace-driven/ISA_Def/volta_opcode.h b/src/trace-driven/ISA_Def/volta_opcode.h new file mode 100644 index 0000000..03d50b7 --- /dev/null +++ b/src/trace-driven/ISA_Def/volta_opcode.h @@ -0,0 +1,175 @@ +//developed by Mahmoud Khairy, Purdue Univ +//abdallm@purdue.edu + +#ifndef VOLTA_OPCODE_H +#define VOLTA_OPCODE_H + +#include "trace_opcode.h" +#include +#include + +#define VOLTA_BINART_VERSION 70 +#define VOLTA_SHARED_MEMORY_VIRTIAL_ADDRESS_START 0x00007f2c60000000 + +//TO DO: moving this to a yml or def files + +///Volta SM_70 ISA +//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html +static const std::unordered_map Volta_OpcodeMap = { + //Floating Point 32 Instructions + {"FADD", OpcodeChar(OP_FADD, SP_OP)}, + {"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)}, + {"FCHK", OpcodeChar(OP_FCHK, SP_OP)}, + {"FFMA32I", OpcodeChar(OP_FFMA32I, SP_OP)}, + {"FFMA", OpcodeChar(OP_FFMA, SP_OP)}, + {"FMNMX", OpcodeChar(OP_FMNMX, SP_OP)}, + {"FMUL", OpcodeChar(OP_FMUL, SP_OP)}, + {"FMUL32I", OpcodeChar(OP_FMUL32I, SP_OP)}, + {"FSEL", OpcodeChar(OP_FSEL, SP_OP)}, + {"FSET", OpcodeChar(OP_FSET, SP_OP)}, + {"FSETP", OpcodeChar(OP_FSETP, SP_OP)}, + {"FSWZADD", OpcodeChar(OP_FSWZADD, SP_OP)}, + //SFU + {"MUFU", OpcodeChar(OP_MUFU, SFU_OP)}, + + //Floating Point 16 Instructions + {"HADD2", OpcodeChar(OP_HADD2, SP_OP)}, + {"HADD2_32I", OpcodeChar(OP_HADD2_32I, SP_OP)}, + {"HFMA2", OpcodeChar(OP_HFMA2, SP_OP)}, + {"HFMA2_32I", OpcodeChar(OP_HFMA2_32I, SP_OP)}, + {"HMUL2", OpcodeChar(OP_HMUL2, SP_OP)}, + {"HMUL2_32I", OpcodeChar(OP_HMUL2_32I, SP_OP)}, + {"HSET2", OpcodeChar(OP_HSET2, SP_OP)}, + {"HSETP2", OpcodeChar(OP_HSETP2, SP_OP)}, + + //Tensor Core Instructions + {"HMMA", OpcodeChar(OP_HMMA, TENSOR_CORE_OP)}, + + //Double Point Instructions + {"DADD", OpcodeChar(OP_DADD, DP_OP)}, + {"DFMA", OpcodeChar(OP_DFMA, DP_OP)}, + {"DMUL", OpcodeChar(OP_DMUL, DP_OP)}, + {"DSETP", OpcodeChar(OP_DSETP, DP_OP)}, + + //Integer Instructions + {"BMSK", OpcodeChar(OP_BMSK, INTP_OP)}, + {"BREV", OpcodeChar(OP_BREV, INTP_OP)}, + {"FLO", OpcodeChar(OP_FLO, INTP_OP)}, + {"IABS", OpcodeChar(OP_IABS, INTP_OP)}, + {"IADD", OpcodeChar(OP_IADD, INTP_OP)}, + {"IADD3", OpcodeChar(OP_IADD3, INTP_OP)}, + {"IADD32I", OpcodeChar(OP_IADD32I, INTP_OP)}, + {"IDP", OpcodeChar(OP_IDP, INTP_OP)}, + {"IDP4A", OpcodeChar(OP_IDP4A, INTP_OP)}, + {"IMAD", OpcodeChar(OP_IMAD, INTP_OP)}, + {"IMMA", OpcodeChar(OP_IMMA, INTP_OP)}, + {"IMNMX", OpcodeChar(OP_IMNMX, INTP_OP)}, + {"IMUL", OpcodeChar(OP_IMUL, INTP_OP)}, + {"IMUL32I", OpcodeChar(OP_IMUL32I, INTP_OP)}, + {"ISCADD", OpcodeChar(OP_ISCADD, INTP_OP)}, + {"ISCADD32I", OpcodeChar(OP_ISCADD32I, INTP_OP)}, + {"ISETP", OpcodeChar(OP_ISETP, INTP_OP)}, + {"LEA", OpcodeChar(OP_LEA, INTP_OP)}, + {"LOP", OpcodeChar(OP_LOP, INTP_OP)}, + {"LOP3", OpcodeChar(OP_LOP3, INTP_OP)}, + {"LOP32I", OpcodeChar(OP_LOP32I, INTP_OP)}, + {"POPC", OpcodeChar(OP_POPC, INTP_OP)}, + {"SHF", OpcodeChar(OP_SHF, INTP_OP)}, + {"SHR", OpcodeChar(OP_SHR, INTP_OP)}, + {"VABSDIFF", OpcodeChar(OP_VABSDIFF, INTP_OP)}, + {"VABSDIFF4", OpcodeChar(OP_VABSDIFF4, INTP_OP)}, + + //Conversion Instructions + {"F2F", OpcodeChar(OP_F2F, ALU_OP)}, + {"F2I", OpcodeChar(OP_F2I, ALU_OP)}, + {"I2F", OpcodeChar(OP_I2F, ALU_OP)}, + {"I2I", OpcodeChar(OP_I2I, ALU_OP)}, + {"I2IP", OpcodeChar(OP_I2IP, ALU_OP)}, + {"FRND", OpcodeChar(OP_FRND, ALU_OP)}, + + //Movement Instructions + {"MOV", OpcodeChar(OP_MOV, ALU_OP)}, + {"MOV32I", OpcodeChar(OP_MOV32I, ALU_OP)}, + {"PRMT", OpcodeChar(OP_PRMT, ALU_OP)}, + {"SEL", OpcodeChar(OP_SEL, ALU_OP)}, + {"SGXT", OpcodeChar(OP_SGXT, ALU_OP)}, + {"SHFL", OpcodeChar(OP_SHFL, ALU_OP)}, + + //Predicate Instructions + {"PLOP3", OpcodeChar(OP_PLOP3, ALU_OP)}, + {"PSETP", OpcodeChar(OP_PSETP, ALU_OP)}, + {"P2R", OpcodeChar(OP_P2R, ALU_OP)}, + {"R2P", OpcodeChar(OP_R2P, ALU_OP)}, + + //Load/Store Instructions + {"LD", OpcodeChar(OP_LD, LOAD_OP)}, + //For now, we ignore constant loads, consider it as ALU_OP, TO DO + {"LDC", OpcodeChar(OP_LDC, ALU_OP)}, + {"LDG", OpcodeChar(OP_LDG, LOAD_OP)}, + {"LDL", OpcodeChar(OP_LDL, LOAD_OP)}, + {"LDS", OpcodeChar(OP_LDS, LOAD_OP)}, + {"ST", OpcodeChar(OP_ST, STORE_OP)}, + {"STG", OpcodeChar(OP_STG, STORE_OP)}, + {"STL", OpcodeChar(OP_STL, STORE_OP)}, + {"STS", OpcodeChar(OP_STS, STORE_OP)}, + {"MATCH", OpcodeChar(OP_MATCH, ALU_OP)}, + {"QSPC", OpcodeChar(OP_QSPC, ALU_OP)}, + {"ATOM", OpcodeChar(OP_ATOM, STORE_OP)}, + {"ATOMS", OpcodeChar(OP_ATOMS, STORE_OP)}, + {"ATOMG", OpcodeChar(OP_ATOMG, STORE_OP)}, + {"RED", OpcodeChar(OP_RED, STORE_OP)}, + {"CCTL", OpcodeChar(OP_CCTL, ALU_OP)}, + {"CCTLL", OpcodeChar(OP_CCTLL, ALU_OP)}, + {"ERRBAR", OpcodeChar(OP_ERRBAR, ALU_OP)}, + {"MEMBAR", OpcodeChar(OP_MEMBAR, MEMORY_BARRIER_OP)}, + {"CCTLT", OpcodeChar(OP_CCTLT, ALU_OP)}, + + //Texture Instructions + //For now, we ignore texture loads, consider it as ALU_OP + {"TEX", OpcodeChar(OP_TEX, ALU_OP)}, + {"TLD", OpcodeChar(OP_TLD, ALU_OP)}, + {"TLD4", OpcodeChar(OP_TLD4, ALU_OP)}, + {"TMML", OpcodeChar(OP_TMML, ALU_OP)}, + {"TXD", OpcodeChar(OP_TXD, ALU_OP)}, + {"TXQ", OpcodeChar(OP_TXQ, ALU_OP)}, + + //Control Instructions + {"BMOV", OpcodeChar(OP_BMOV, BRANCH_OP)}, + {"BPT", OpcodeChar(OP_BPT, BRANCH_OP)}, + {"BRA", OpcodeChar(OP_BRA, BRANCH_OP)}, + {"BREAK", OpcodeChar(OP_BREAK, BRANCH_OP)}, + {"BRX", OpcodeChar(OP_BRX, BRANCH_OP)}, + {"BSSY", OpcodeChar(OP_BSSY, BRANCH_OP)}, + {"BSYNC", OpcodeChar(OP_BSYNC, BRANCH_OP)}, + {"CALL", OpcodeChar(OP_CALL, CALL_OPS)}, + {"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)}, + {"JMP", OpcodeChar(OP_JMP, BRANCH_OP)}, + {"JMX", OpcodeChar(OP_JMX, BRANCH_OP)}, + {"KILL", OpcodeChar(OP_KILL, BRANCH_OP)}, + {"NANOSLEEP", OpcodeChar(OP_NANOSLEEP, BRANCH_OP)}, + {"RET", OpcodeChar(OP_RET, RET_OPS)}, + {"RPCMOV", OpcodeChar(OP_RPCMOV, BRANCH_OP)}, + {"RTT", OpcodeChar(OP_RTT, RET_OPS)}, + {"WARPSYNC", OpcodeChar(OP_WARPSYNC, BRANCH_OP)}, + {"YIELD", OpcodeChar(OP_YIELD, BRANCH_OP)}, + + //Miscellaneous Instructions + {"B2R", OpcodeChar(OP_B2R, ALU_OP)}, + {"BAR", OpcodeChar(OP_BAR, BARRIER_OP)}, + {"CS2R", OpcodeChar(OP_CS2R, ALU_OP)}, + {"CSMTEST", OpcodeChar(OP_CSMTEST, ALU_OP)}, + {"DEPBAR", OpcodeChar(OP_DEPBAR, ALU_OP)}, + {"GETLMEMBASE", OpcodeChar(OP_GETLMEMBASE, ALU_OP)}, + {"LEPC", OpcodeChar(OP_LEPC ,ALU_OP)}, + {"NOP", OpcodeChar(OP_NOP ,ALU_OP)}, + {"PMTRIG", OpcodeChar(OP_PMTRIG, ALU_OP)}, + {"R2B", OpcodeChar(OP_R2B, ALU_OP)}, + {"S2R", OpcodeChar(OP_S2R, ALU_OP)}, + {"SETCTAID", OpcodeChar(OP_SETCTAID, ALU_OP)}, + {"SETLMEMBASE", OpcodeChar(OP_SETLMEMBASE, ALU_OP)}, + {"VOTE", OpcodeChar(OP_VOTE, ALU_OP)}, + {"VOTE_VTG", OpcodeChar(OP_VOTE_VTG, ALU_OP)}, + +}; + +#endif diff --git a/src/trace-driven/gpgpusim_trace_driven_main.cc b/src/trace-driven/gpgpusim_trace_driven_main.cc index 5e07ace..90dc769 100644 --- a/src/trace-driven/gpgpusim_trace_driven_main.cc +++ b/src/trace-driven/gpgpusim_trace_driven_main.cc @@ -16,7 +16,7 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../../libcuda/gpgpu_context.h" #include "trace_driven.h" -#include "trace_opcode.h" +#include "ISA_Def/trace_opcode.h" #include "../gpgpusim_entrypoint.h" /* TO DO: diff --git a/src/trace-driven/kepler_opcode.h b/src/trace-driven/kepler_opcode.h deleted file mode 100644 index 4aa8e0f..0000000 --- a/src/trace-driven/kepler_opcode.h +++ /dev/null @@ -1,150 +0,0 @@ -//developed by Mahmoud Khairy, Purdue Univ -//abdallm@purdue.edu - -#ifndef KEPLER_OPCODE_H -#define KEPLER_OPCODE_H - -#include "../abstract_hardware_model.h" -#include "trace_opcode.h" -#include -#include - -#define KEPLER_BINART_VERSION 35 -#define KEPLER_SHARED_MEMORY_VIRTIAL_ADDRESS_START 0x00007f2c60000000 - -//TO DO: moving this to a yml or def files - -///Kepler ISA -//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html -static const std::unordered_map Kepler_OpcodeMap = { - //Floating Point 32 Instructions - {"FFMA", OpcodeChar(OP_FFMA, SP_OP)}, - {"FFMA32I", OpcodeChar(OP_FFMA32I, SP_OP)}, - {"FADD", OpcodeChar(OP_FADD, SP_OP)}, - {"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)}, - {"FCMP", OpcodeChar(OP_FCMP, SP_OP)}, - {"FMUL", OpcodeChar(OP_FMUL, SP_OP)}, - {"FMUL32I", OpcodeChar(OP_FMUL32I, SP_OP)}, - {"FMNMX", OpcodeChar(OP_FMNMX, SP_OP)}, - {"FSWZ", OpcodeChar(OP_FSWZ, SP_OP)}, - {"FSET", OpcodeChar(OP_FSET, SP_OP)}, - {"FSETP", OpcodeChar(OP_FSETP, SP_OP)}, - {"FCHK", OpcodeChar(OP_FCHK, SP_OP)}, - {"RRO", OpcodeChar(OP_RRO, SP_OP)}, - //SFU - {"MUFU", OpcodeChar(OP_MUFU, SFU_OP)}, - - - //Double Point Instructions - {"DFMA", OpcodeChar(OP_DFMA, DP_OP)}, - {"DADD", OpcodeChar(OP_DADD, DP_OP)}, - {"DMUL", OpcodeChar(OP_DMUL, DP_OP)}, - {"DMNMX", OpcodeChar(OP_DMNMX, DP_OP)}, - {"DSET", OpcodeChar(OP_DSET, DP_OP)}, - {"DSETP", OpcodeChar(OP_DSETP, DP_OP)}, - - //Integer Instructions - {"IMAD", OpcodeChar(OP_IMAD, INTP_OP)}, - {"IMADSP", OpcodeChar(OP_IMADSP, INTP_OP)}, - {"IMUL", OpcodeChar(OP_IMUL, INTP_OP)}, - {"IMUL32I", OpcodeChar(OP_IMUL32I, INTP_OP)}, - {"IADD", OpcodeChar(OP_IADD, INTP_OP)}, - {"IADD32I", OpcodeChar(OP_IADD32I, INTP_OP)}, - {"ISUB", OpcodeChar(OP_ISUB, INTP_OP)}, - {"ISCADD", OpcodeChar(OP_ISCADD, INTP_OP)}, - {"ISCADD32I", OpcodeChar(OP_ISCADD32I, INTP_OP)}, - {"ISAD", OpcodeChar(OP_ISAD, INTP_OP)}, - {"IMNMX", OpcodeChar(OP_IMNMX, INTP_OP)}, - {"BFE", OpcodeChar(OP_BFE, INTP_OP)}, - {"BFI", OpcodeChar(OP_BFI, INTP_OP)}, - {"SHR", OpcodeChar(OP_SHR, INTP_OP)}, - {"SHL", OpcodeChar(OP_SHL, INTP_OP)}, - {"SHF", OpcodeChar(OP_SHF, INTP_OP)}, - {"LOP", OpcodeChar(OP_LOP, INTP_OP)}, - {"LOP32I", OpcodeChar(OP_LOP32I, INTP_OP)}, - {"FLO", OpcodeChar(OP_FLO, INTP_OP)}, - {"ISET", OpcodeChar(OP_ISET, INTP_OP)}, - {"ISETP", OpcodeChar(OP_ISETP, INTP_OP)}, - {"ICMP", OpcodeChar(OP_ICMP, INTP_OP)}, - {"POPC", OpcodeChar(OP_POPC, INTP_OP)}, - - //Conversion Instructions - {"F2F", OpcodeChar(OP_F2F, ALU_OP)}, - {"F2I", OpcodeChar(OP_F2I, ALU_OP)}, - {"I2F", OpcodeChar(OP_I2F, ALU_OP)}, - {"I2I", OpcodeChar(OP_I2I, ALU_OP)}, - - //Movement Instructions - {"MOV", OpcodeChar(OP_MOV, ALU_OP)}, - {"MOV32I", OpcodeChar(OP_MOV32I, ALU_OP)}, - {"SEL", OpcodeChar(OP_SEL, ALU_OP)}, - {"PRMT", OpcodeChar(OP_PRMT, ALU_OP)}, - {"SHFL", OpcodeChar(OP_SHFL, ALU_OP)}, - - //Predicate Instructions - {"P2R", OpcodeChar(OP_P2R, ALU_OP)}, - {"R2P", OpcodeChar(OP_R2P, ALU_OP)}, - {"CSET", OpcodeChar(OP_CSET, ALU_OP)}, - {"CSETP", OpcodeChar(OP_CSETP, ALU_OP)}, - {"PSET", OpcodeChar(OP_PSET, ALU_OP)}, - {"PSETP", OpcodeChar(OP_PSETP, ALU_OP)}, - - //Texture Instructions - //For now, we ignore texture loads, consider it as ALU_OP - {"TEX", OpcodeChar(OP_TEX, ALU_OP)}, - {"TLD", OpcodeChar(OP_TLD, ALU_OP)}, - {"TLD4", OpcodeChar(OP_TLD4, ALU_OP)}, - {"TXQ", OpcodeChar(OP_TXQ, ALU_OP)}, - - //Load/Store Instructions - //For now, we ignore constant loads, consider it as ALU_OP, TO DO - {"LDC", OpcodeChar(OP_LDC, ALU_OP)}, - //in Kepler, LD is load global so set it to LDG - {"LD", OpcodeChar(OP_LDG, LOAD_OP)}, - {"LDG", OpcodeChar(OP_LDG, LOAD_OP)}, - {"LDL", OpcodeChar(OP_LDL, LOAD_OP)}, - {"LDS", OpcodeChar(OP_LDS, LOAD_OP)}, - {"LDSLK", OpcodeChar(OP_LDSLK, LOAD_OP)}, - {"ST", OpcodeChar(OP_STG, STORE_OP)}, - {"STL", OpcodeChar(OP_STL, STORE_OP)}, - {"STS", OpcodeChar(OP_STS, STORE_OP)}, - {"STSCUL", OpcodeChar(OP_STSCUL, STORE_OP)}, - {"ATOM", OpcodeChar(OP_ATOM, STORE_OP)}, - {"RED", OpcodeChar(OP_RED, STORE_OP)}, - {"CCTL", OpcodeChar(OP_CCTL, ALU_OP)}, - {"CCTLL", OpcodeChar(OP_CCTLL, ALU_OP)}, - {"MEMBAR", OpcodeChar(OP_MEMBAR, MEMORY_BARRIER_OP)}, - - //surface memory instructions - {"SUCLAMP", OpcodeChar(OP_SUCLAMP, LOAD_OP)}, - {"SUBFM", OpcodeChar(OP_SUBFM, LOAD_OP)}, - {"SUEAU", OpcodeChar(OP_SUEAU, LOAD_OP)}, - {"SULDGA", OpcodeChar(OP_SULDGA, LOAD_OP)}, - {"SUSTGA", OpcodeChar(OP_SUSTGA, STORE_OP)}, - - //Control Instructions - {"BRA", OpcodeChar(OP_BRA, BRANCH_OP)}, - {"BRX", OpcodeChar(OP_BRX, BRANCH_OP)}, - {"JMP", OpcodeChar(OP_JMP, BRANCH_OP)}, - {"JMX", OpcodeChar(OP_JMX, BRANCH_OP)}, - {"CAL", OpcodeChar(OP_CAL, CALL_OPS)}, - {"JCAL", OpcodeChar(OP_JCAL, CALL_OPS)}, - {"RET", OpcodeChar(OP_RET, RET_OPS)}, - {"BRK", OpcodeChar(OP_BRK, RET_OPS)}, - {"CONT", OpcodeChar(OP_CONT, RET_OPS)}, - {"SSY", OpcodeChar(OP_SSY, RET_OPS)}, - {"PBK", OpcodeChar(OP_PBK, RET_OPS)}, - {"PCNT", OpcodeChar(OP_PCNT, RET_OPS)}, - {"PRET", OpcodeChar(OP_PRET, RET_OPS)}, - {"BPT", OpcodeChar(OP_BPT, BRANCH_OP)}, - {"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)}, - - //Miscellaneous Instructions - {"NOP", OpcodeChar(OP_NOP, ALU_OP)}, - {"S2R", OpcodeChar(OP_S2R, ALU_OP)}, - {"B2R", OpcodeChar(OP_B2R, ALU_OP)}, - {"BAR", OpcodeChar(OP_BAR, BARRIER_OP)}, - {"VOTE", OpcodeChar(OP_VOTE, ALU_OP)}, -}; - -#endif diff --git a/src/trace-driven/pascal_opcode.h b/src/trace-driven/pascal_opcode.h deleted file mode 100644 index 2cacb28..0000000 --- a/src/trace-driven/pascal_opcode.h +++ /dev/null @@ -1,202 +0,0 @@ -//developed by Mahmoud Khairy, Purdue Univ -//abdallm@purdue.edu - -#ifndef PASCAL_OPCODE_H -#define PASCAL_OPCODE_H - -#include "../abstract_hardware_model.h" -#include "trace_opcode.h" -#include -#include - -#define PASCAL_TITANX_BINART_VERSION 61 -#define PASCAL_P100_BINART_VERSION 60 - -#define PASCAL_SHARED_MEMORY_VIRTIAL_ADDRESS_START 0x00007f2c60000000 - -//TO DO: moving this to a yml or def files - -///Pascal SM_61 ISA -//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html -static const std::unordered_map Pascal_OpcodeMap = { - //Floating Point 32 Instructions - {"FADD", OpcodeChar(OP_FADD, SP_OP)}, - {"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)}, - {"FCHK", OpcodeChar(OP_FCHK, SP_OP)}, - {"FFMA32I", OpcodeChar(OP_FFMA32I, SP_OP)}, - {"FFMA", OpcodeChar(OP_FFMA, SP_OP)}, - {"FMNMX", OpcodeChar(OP_FMNMX, SP_OP)}, - {"FMUL", OpcodeChar(OP_FMUL, SP_OP)}, - {"FMUL32I", OpcodeChar(OP_FMUL32I, SP_OP)}, - {"FSEL", OpcodeChar(OP_FSEL, SP_OP)}, - {"FSET", OpcodeChar(OP_FSET, SP_OP)}, - {"FSETP", OpcodeChar(OP_FSETP, SP_OP)}, - {"FSWZADD", OpcodeChar(OP_FSWZADD, SP_OP)}, - {"RRO", OpcodeChar(OP_RRO, SP_OP)}, - - //SFU - {"MUFU", OpcodeChar(OP_MUFU, SFU_OP)}, - - //Floating Point 16 Instructions - {"HADD2", OpcodeChar(OP_HADD2, SP_OP)}, - {"HFMA2", OpcodeChar(OP_HFMA2, SP_OP)}, - {"HMUL2", OpcodeChar(OP_HMUL2, SP_OP)}, - {"HSET2", OpcodeChar(OP_HSET2, SP_OP)}, - {"HSETP2", OpcodeChar(OP_HSETP2, SP_OP)}, - - //Double Point Instructions - {"DADD", OpcodeChar(OP_DADD, DP_OP)}, - {"DFMA", OpcodeChar(OP_DFMA, DP_OP)}, - {"DMUL", OpcodeChar(OP_DMUL, DP_OP)}, - {"DSETP", OpcodeChar(OP_DSETP, DP_OP)}, - {"DMNMX", OpcodeChar(OP_DMNMX, DP_OP)}, - {"DSET", OpcodeChar(OP_DSET, DP_OP)}, - - //Integer Instructions - {"BMSK", OpcodeChar(OP_BMSK, INTP_OP)}, - {"BREV", OpcodeChar(OP_BREV, INTP_OP)}, - {"FLO", OpcodeChar(OP_FLO, INTP_OP)}, - {"IABS", OpcodeChar(OP_IABS, INTP_OP)}, - {"IADD", OpcodeChar(OP_IADD, INTP_OP)}, - {"IADD3", OpcodeChar(OP_IADD3, INTP_OP)}, - {"IADD32I", OpcodeChar(OP_IADD32I, INTP_OP)}, - {"IDP", OpcodeChar(OP_IDP, INTP_OP)}, - {"IDP4A", OpcodeChar(OP_IDP4A, INTP_OP)}, - {"IMAD", OpcodeChar(OP_IMAD, INTP_OP)}, - {"IMMA", OpcodeChar(OP_IMMA, INTP_OP)}, - {"IMNMX", OpcodeChar(OP_IMNMX, INTP_OP)}, - {"IMUL", OpcodeChar(OP_IMUL, INTP_OP)}, - {"IMUL32I", OpcodeChar(OP_IMUL32I, INTP_OP)}, - {"ISCADD", OpcodeChar(OP_ISCADD, INTP_OP)}, - {"ISCADD32I", OpcodeChar(OP_ISCADD32I, INTP_OP)}, - {"ISETP", OpcodeChar(OP_ISETP, INTP_OP)}, - {"ISET", OpcodeChar(OP_ISET, INTP_OP)}, - {"LEA", OpcodeChar(OP_LEA, INTP_OP)}, - {"LOP", OpcodeChar(OP_LOP, INTP_OP)}, - {"LOP3", OpcodeChar(OP_LOP3, INTP_OP)}, - {"LOP32I", OpcodeChar(OP_LOP32I, INTP_OP)}, - {"POPC", OpcodeChar(OP_POPC, INTP_OP)}, - {"SHF", OpcodeChar(OP_SHF, INTP_OP)}, - {"SHR", OpcodeChar(OP_SHR, INTP_OP)}, - {"VABSDIFF", OpcodeChar(OP_VABSDIFF, INTP_OP)}, - {"VABSDIFF4", OpcodeChar(OP_VABSDIFF4, INTP_OP)}, - {"BFE", OpcodeChar(OP_BFE, INTP_OP)}, - {"BFI", OpcodeChar(OP_BFI, INTP_OP)}, - {"ICMP", OpcodeChar(OP_ICMP, INTP_OP)}, - {"IMADSP", OpcodeChar(OP_IMADSP, INTP_OP)}, - {"SHL", OpcodeChar(OP_SHL, INTP_OP)}, - {"XMAD", OpcodeChar(OP_XMAD, INTP_OP)}, - {"VMNMX", OpcodeChar(OP_VMNMX, INTP_OP)}, - - - //Conversion Instructions - {"F2F", OpcodeChar(OP_F2F, ALU_OP)}, - {"F2I", OpcodeChar(OP_F2I, ALU_OP)}, - {"I2F", OpcodeChar(OP_I2F, ALU_OP)}, - {"I2I", OpcodeChar(OP_I2I, ALU_OP)}, - {"I2IP", OpcodeChar(OP_I2IP, ALU_OP)}, - {"FRND", OpcodeChar(OP_FRND, ALU_OP)}, - - //Movement Instructions - {"MOV", OpcodeChar(OP_MOV, ALU_OP)}, - {"MOV32I", OpcodeChar(OP_MOV32I, ALU_OP)}, - {"PRMT", OpcodeChar(OP_PRMT, ALU_OP)}, - {"SEL", OpcodeChar(OP_SEL, ALU_OP)}, - {"SGXT", OpcodeChar(OP_SGXT, ALU_OP)}, - {"SHFL", OpcodeChar(OP_SHFL, ALU_OP)}, - - //Predicate Instructions - {"PLOP3", OpcodeChar(OP_PLOP3, ALU_OP)}, - {"PSETP", OpcodeChar(OP_PSETP, ALU_OP)}, - {"P2R", OpcodeChar(OP_P2R, ALU_OP)}, - {"R2P", OpcodeChar(OP_R2P, ALU_OP)}, - {"CSET", OpcodeChar(OP_CSET, ALU_OP)}, - {"CSETP", OpcodeChar(OP_CSETP, ALU_OP)}, - {"PSET", OpcodeChar(OP_PSET, ALU_OP)}, - - - //Load/Store Instructions - {"LD", OpcodeChar(OP_LD, LOAD_OP)}, - //For now, we ignore constant loads, consider it as ALU_OP, TO DO - {"LDC", OpcodeChar(OP_LDC, ALU_OP)}, - {"LDG", OpcodeChar(OP_LDG, LOAD_OP)}, - {"LDL", OpcodeChar(OP_LDL, LOAD_OP)}, - {"LDS", OpcodeChar(OP_LDS, LOAD_OP)}, - {"ST", OpcodeChar(OP_ST, STORE_OP)}, - {"STG", OpcodeChar(OP_STG, STORE_OP)}, - {"STL", OpcodeChar(OP_STL, STORE_OP)}, - {"STS", OpcodeChar(OP_STS, STORE_OP)}, - {"MATCH", OpcodeChar(OP_MATCH, ALU_OP)}, - {"QSPC", OpcodeChar(OP_QSPC, ALU_OP)}, - {"ATOM", OpcodeChar(OP_ATOM, STORE_OP)}, - {"ATOMS", OpcodeChar(OP_ATOMS, STORE_OP)}, - {"ATOMG", OpcodeChar(OP_ATOMG, STORE_OP)}, - {"RED", OpcodeChar(OP_RED, STORE_OP)}, - {"CCTL", OpcodeChar(OP_CCTL, ALU_OP)}, - {"CCTLL", OpcodeChar(OP_CCTLL, ALU_OP)}, - {"ERRBAR", OpcodeChar(OP_ERRBAR, ALU_OP)}, - {"MEMBAR", OpcodeChar(OP_MEMBAR, MEMORY_BARRIER_OP)}, - {"CCTLT", OpcodeChar(OP_CCTLT, ALU_OP)}, - - //Texture Instructions - //For now, we ignore texture loads, consider it as ALU_OP - {"TEX", OpcodeChar(OP_TEX, ALU_OP)}, - {"TLD", OpcodeChar(OP_TLD, ALU_OP)}, - {"TLD4", OpcodeChar(OP_TLD4, ALU_OP)}, - {"TMML", OpcodeChar(OP_TMML, ALU_OP)}, - {"TXD", OpcodeChar(OP_TXD, ALU_OP)}, - {"TXQ", OpcodeChar(OP_TXQ, ALU_OP)}, - {"TEXS", OpcodeChar(OP_TEXS, ALU_OP)}, - {"TLD4S", OpcodeChar(OP_TLD4S, ALU_OP)}, - {"TLDS", OpcodeChar(OP_TLDS, ALU_OP)}, - - //Control Instructions - {"BMOV", OpcodeChar(OP_BMOV, BRANCH_OP)}, - {"BPT", OpcodeChar(OP_BPT, BRANCH_OP)}, - {"BRA", OpcodeChar(OP_BRA, BRANCH_OP)}, - {"BREAK", OpcodeChar(OP_BREAK, BRANCH_OP)}, - {"BRX", OpcodeChar(OP_BRX, BRANCH_OP)}, - {"BSSY", OpcodeChar(OP_BSSY, BRANCH_OP)}, - {"BSYNC", OpcodeChar(OP_BSYNC, BRANCH_OP)}, - {"CALL", OpcodeChar(OP_CALL, CALL_OPS)}, - {"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)}, - {"JMP", OpcodeChar(OP_JMP, BRANCH_OP)}, - {"SSY", OpcodeChar(OP_SSY, BRANCH_OP)}, - {"SYNC", OpcodeChar(OP_SYNC, BRANCH_OP)}, - {"JMX", OpcodeChar(OP_JMX, BRANCH_OP)}, - {"KILL", OpcodeChar(OP_KILL, BRANCH_OP)}, - {"NANOSLEEP", OpcodeChar(OP_NANOSLEEP, BRANCH_OP)}, - {"RET", OpcodeChar(OP_RET, RET_OPS)}, - {"RPCMOV", OpcodeChar(OP_RPCMOV, BRANCH_OP)}, - {"RTT", OpcodeChar(OP_RTT, RET_OPS)}, - {"WARPSYNC", OpcodeChar(OP_WARPSYNC, BRANCH_OP)}, - {"YIELD", OpcodeChar(OP_YIELD, BRANCH_OP)}, - {"CAL", OpcodeChar(OP_CAL, CALL_OPS)}, - {"JCAL", OpcodeChar(OP_JCAL, CALL_OPS)}, - {"PRET", OpcodeChar(OP_PRET, CALL_OPS)}, - {"BRK", OpcodeChar(OP_BRK, CALL_OPS)}, - {"PBK", OpcodeChar(OP_PBK, CALL_OPS)}, - {"CONT", OpcodeChar(OP_CONT, CALL_OPS)}, - {"PCNT", OpcodeChar(OP_PCNT, CALL_OPS)}, - {"PEXIT", OpcodeChar(OP_PEXIT, CALL_OPS)}, - - //Miscellaneous Instructions - {"B2R", OpcodeChar(OP_B2R, ALU_OP)}, - {"BAR", OpcodeChar(OP_BAR, BARRIER_OP)}, - {"CS2R", OpcodeChar(OP_CS2R, ALU_OP)}, - {"CSMTEST", OpcodeChar(OP_CSMTEST, ALU_OP)}, - {"DEPBAR", OpcodeChar(OP_DEPBAR, ALU_OP)}, - {"GETLMEMBASE", OpcodeChar(OP_GETLMEMBASE, ALU_OP)}, - {"LEPC", OpcodeChar(OP_LEPC ,ALU_OP)}, - {"NOP", OpcodeChar(OP_NOP ,ALU_OP)}, - {"PMTRIG", OpcodeChar(OP_PMTRIG, ALU_OP)}, - {"R2B", OpcodeChar(OP_R2B, ALU_OP)}, - {"S2R", OpcodeChar(OP_S2R, ALU_OP)}, - {"SETCTAID", OpcodeChar(OP_SETCTAID, ALU_OP)}, - {"SETLMEMBASE", OpcodeChar(OP_SETLMEMBASE, ALU_OP)}, - {"VOTE", OpcodeChar(OP_VOTE, ALU_OP)}, - {"VOTE_VTG", OpcodeChar(OP_VOTE_VTG, ALU_OP)}, - -}; - -#endif diff --git a/src/trace-driven/trace_driven.cc b/src/trace-driven/trace_driven.cc index 22c527e..76eb7ca 100644 --- a/src/trace-driven/trace_driven.cc +++ b/src/trace-driven/trace_driven.cc @@ -19,11 +19,11 @@ #include "../gpgpu-sim/gpu-sim.h" #include "../../libcuda/gpgpu_context.h" #include "trace_driven.h" -#include "trace_opcode.h" -#include "volta_opcode.h" -#include "turing_opcode.h" -#include "pascal_opcode.h" -#include "kepler_opcode.h" +#include "ISA_Def/trace_opcode.h" +#include "ISA_Def/volta_opcode.h" +#include "ISA_Def/turing_opcode.h" +#include "ISA_Def/pascal_opcode.h" +#include "ISA_Def/kepler_opcode.h" #include "../gpgpusim_entrypoint.h" diff --git a/src/trace-driven/trace_driven.h b/src/trace-driven/trace_driven.h index 2888f86..e9fecd9 100644 --- a/src/trace-driven/trace_driven.h +++ b/src/trace-driven/trace_driven.h @@ -9,7 +9,7 @@ #include "../abstract_hardware_model.h" #include "../gpgpu-sim/shader.h" -#include "trace_opcode.h" +#include "ISA_Def/trace_opcode.h" class trace_function_info: public function_info { public: diff --git a/src/trace-driven/trace_opcode.h b/src/trace-driven/trace_opcode.h deleted file mode 100644 index 3492fd3..0000000 --- a/src/trace-driven/trace_opcode.h +++ /dev/null @@ -1,71 +0,0 @@ -//developed by Mahmoud Khairy, Purdue Univ -//abdallm@purdue.edu - -#ifndef TRACE_OPCODE_H -#define TRACE_OPCODE_H - -#include "../abstract_hardware_model.h" -#include -#include - - -enum TraceInstrOpcode { - //volta (common insts for others cards as well) - OP_FADD = 1, OP_FADD32I, OP_FCHK, OP_FFMA32I, OP_FFMA, OP_FMNMX, OP_FMUL, OP_FMUL32I, OP_FSEL, OP_FSET, OP_FSETP, - OP_FSWZADD, OP_MUFU, OP_HADD2, OP_HADD2_32I, OP_HFMA2, OP_HFMA2_32I, OP_HMUL2, OP_HMUL2_32I, OP_HSET2, OP_HSETP2, - OP_HMMA, OP_DADD, OP_DFMA, OP_DMUL, OP_DSETP, - OP_BMSK, OP_BREV, OP_FLO, OP_IABS, OP_IADD, OP_IADD3, OP_IADD32I, OP_IDP, OP_IDP4A, OP_IMAD, OP_IMMA, OP_IMNMX, - OP_IMUL, OP_IMUL32I, OP_ISCADD, OP_ISCADD32I, OP_ISETP, OP_LEA, OP_LOP, OP_LOP3, OP_LOP32I, OP_POPC, OP_SHF, OP_SHR, - OP_VABSDIFF, OP_VABSDIFF4, - OP_F2F, OP_F2I, OP_I2F, OP_I2I, OP_I2IP, OP_FRND, OP_MOV, OP_MOV32I, OP_PRMT, OP_SEL, OP_SGXT, OP_SHFL, OP_PLOP3, - OP_PSETP, OP_P2R, OP_R2P, OP_LD, OP_LDC, OP_LDG, OP_LDL, OP_LDS, OP_ST, OP_STG, OP_STL, OP_STS, OP_MATCH, OP_QSPC, - OP_ATOM, OP_ATOMS, OP_ATOMG, OP_RED, OP_CCTL, OP_CCTLL, OP_ERRBAR, OP_MEMBAR, OP_CCTLT, - OP_TEX, OP_TLD, OP_TLD4, - OP_TMML, OP_TXD, OP_TXQ, OP_BMOV, OP_BPT, OP_BRA, OP_BREAK, OP_BRX, OP_BSSY, OP_BSYNC, OP_CALL, OP_EXIT, OP_JMP, OP_JMX, - OP_KILL, OP_NANOSLEEP, OP_RET, OP_RPCMOV, OP_RTT, OP_WARPSYNC, OP_YIELD, OP_B2R, OP_BAR, OP_CS2R, OP_CSMTEST, OP_DEPBAR, - OP_GETLMEMBASE, OP_LEPC, OP_NOP, OP_PMTRIG, OP_R2B, OP_S2R, OP_SETCTAID, OP_SETLMEMBASE, OP_VOTE, OP_VOTE_VTG, - //unique insts for pascal - OP_RRO, OP_DMNMX, OP_DSET, OP_BFE, OP_BFI, OP_ICMP, OP_IMADSP, OP_SHL, OP_XMAD, OP_CSET, OP_CSETP, - OP_TEXS, OP_TLD4S, OP_TLDS, OP_CAL, OP_JCAL, OP_PRET, OP_BRK, OP_PBK, OP_CONT, OP_PCNT, OP_PEXIT, OP_SSY, OP_SYNC, OP_PSET - , OP_VMNMX, OP_ISET, - //unique insts for kepler - OP_FCMP, OP_FSWZ, OP_ISAD, OP_LDSLK, OP_STSCUL, OP_SUCLAMP, OP_SUBFM, OP_SUEAU, OP_SULDGA, OP_SUSTGA, OP_ISUB, - SASS_NUM_OPCODES /* The total number of opcodes. */ -}; -typedef enum TraceInstrOpcode sass_op_type; - -/* -enum uarch_op_t { - NO_OP=-1, - ALU_OP=1, - SFU_OP, - TENSOR_CORE_OP, - DP_OP, - SP_OP, - INTP_OP, - ALU_SFU_OP, - LOAD_OP, - TENSOR_CORE_LOAD_OP, - TENSOR_CORE_STORE_OP, - STORE_OP, - BRANCH_OP, - BARRIER_OP, - MEMORY_BARRIER_OP, - CALL_OPS, - RET_OPS -}; -typedef enum uarch_op_t op_type; - */ - -struct OpcodeChar -{ - OpcodeChar(unsigned m_opcode, unsigned m_opcode_category) { - opcode = m_opcode; - opcode_category = m_opcode_category; - } - unsigned opcode; - unsigned opcode_category; -}; - - -#endif diff --git a/src/trace-driven/traces-generator/NVbit_tool/Makefile b/src/trace-driven/traces-generator/NVbit_tool/Makefile deleted file mode 100644 index fc0209a..0000000 --- a/src/trace-driven/traces-generator/NVbit_tool/Makefile +++ /dev/null @@ -1,23 +0,0 @@ -NVCC=nvcc -ccbin=`which gcc` -D_FORCE_INLINES -NVBIT_PATH=../../core -INCLUDES=-I$(NVBIT_PATH) -LIBS=-L$(NVBIT_PATH) -lnvbit -NVCC_PATH=-L $(subst bin/nvcc,lib64,$(shell which nvcc | tr -s /)) -SOURCES=$(wildcard *.cu) -OBJECTS=$(SOURCES:.cu=.o) -ARCH=35 - -mkfile_path := $(abspath $(lastword $(MAKEFILE_LIST))) -current_dir := $(notdir $(patsubst %/,%,$(dir $(mkfile_path)))) - -all: $(OBJECTS) $(NVBIT_PATH)/libnvbit.a - $(NVCC) -arch=sm_$(ARCH) -O3 *.o $(LIBS) $(NVCC_PATH) -lcuda -lcudart_static -shared -o ${current_dir}.so - -%.o: %.cu - $(NVCC) -dc -c -std=c++11 $(INCLUDES) -Xptxas -cloning=no -maxrregcount=16 -Xcompiler -Wall -arch=sm_$(ARCH) -O3 -Xcompiler -fPIC $< -o $@ - -$(NVBIT_PATH)/libnvbit.a: - make -C $(NVBIT_PATH) - -clean: - rm -f *.so *.o diff --git a/src/trace-driven/traces-generator/NVbit_tool/traceall.cu b/src/trace-driven/traces-generator/NVbit_tool/traceall.cu deleted file mode 100644 index 1426528..0000000 --- a/src/trace-driven/traces-generator/NVbit_tool/traceall.cu +++ /dev/null @@ -1,646 +0,0 @@ -/* Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved. - * - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions - * are met: - * * Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * * Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution. - * * Neither the name of NVIDIA CORPORATION nor the names of its - * contributors may be used to endorse or promote products derived - * from this software without specific prior written permission. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY - * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE - * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR - * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR - * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, - * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, - * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR - * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY - * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE - * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - */ - -/* Author: Oreste Villa, ovilla@nvidia.com - 2018 */ - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -/* every tool needs to include this once */ -#include "nvbit_tool.h" - -/* nvbit interface file */ -#include "nvbit.h" - -/* for channel */ -#include "utils/channel.hpp" - -/* for _cuda_safe and GET_VAR* macros */ -#include "macros.h" - -/* Channel used to communicate from GPU to CPU receiving thread */ -#define CHANNEL_SIZE (1l << 20) -static __managed__ ChannelDev channel_dev; -static ChannelHost channel_host; - -/* receiving thread and its control variables */ -pthread_t recv_thread; -volatile bool recv_thread_started = false; -volatile bool recv_thread_receiving = false; - -/* skip flag used to avoid re-entry on the nvbit_callback when issuing - * flush_channel kernel call */ -bool skip_flag = false; - -/* global control variables for this tool */ -uint32_t instr_begin_interval = 0; -uint32_t instr_end_interval = UINT32_MAX; -int verbose = 0; - -/* opcode to id map and reverse map */ -std::map opcode_to_id_map; -std::map id_to_opcode_map; - -/* kernel instruction counter, updated by the GPU */ -static __managed__ uint64_t total_dynamic_instr_counter = 0; -static __managed__ uint64_t reported_dynamic_instr_counter = 0; -static __managed__ uint64_t dynamic_instr_limit = 0; -uint64_t dynamic_instr_limit_input = 0; //0 means no limit - -#define MAX_SRC 4 -/* information collected in the instrumentation function */ -typedef struct { - int cta_id_x; - int cta_id_y; - int cta_id_z; - int warpid_tb; - int warpid_sm; - int sm_id; - int opcode_id; - uint64_t addrs[32]; - uint32_t vpc; - bool is_mem; - int32_t GPRDst; - int32_t GPRSrcs[MAX_SRC]; - int32_t numSrcs; - int32_t width; - uint32_t active_mask; - -} mem_access_t; - -/* Instrumentation function that we want to inject, please note the use of - * 1. extern "C" __device__ __noinline__ - * To prevent "dead"-code elimination by the compiler. - * 2. NVBIT_EXPORT_FUNC(dev_func) - * To notify nvbit the name of the function we want to inject. - * This name must match exactly the function name. - */ -extern "C" __device__ __noinline__ void instrument_mem(int pred, int opcode_id, int32_t vpc, - uint32_t reg_high, - uint32_t reg_low, - int32_t imm, - int32_t srcReg1, int32_t srcReg2, int32_t desReg, int32_t width) { - if (!pred) { - return; - } - - uint32_t active_mask = __ballot(1); - const int laneid = get_laneid(); - const int first_laneid = __ffs(active_mask) - 1; - - if (dynamic_instr_limit && total_dynamic_instr_counter >= dynamic_instr_limit) - if (first_laneid == laneid) { - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - return; - } - - mem_access_t ma; - - /* collect memory address information */ - int64_t base_addr = (((uint64_t)reg_high) << 32) | ((uint64_t)reg_low); - uint64_t addr = base_addr + imm; - for (int i = 0; i < 32; i++) { - ma.addrs[i] = __shfl(addr, i); - } - - int4 cta = get_ctaid(); - int uniqe_threadId = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x; - ma.warpid_tb = uniqe_threadId/32; - - ma.cta_id_x = cta.x; - ma.cta_id_y = cta.y; - ma.cta_id_z = cta.z; - ma.warpid_sm = get_warpid(); - ma.opcode_id = opcode_id; - ma.is_mem = true; - ma.vpc = vpc; - ma.width = width; - ma.GPRDst = desReg; - ma.GPRSrcs[0] = srcReg1; - ma.GPRSrcs[1] = srcReg2; - ma.GPRSrcs[2] = -1; - ma.GPRSrcs[3] = -1; - ma.numSrcs = 2; - ma.active_mask = active_mask; - ma.sm_id = get_smid(); - - /* first active lane pushes information on the channel */ - if (first_laneid == laneid) { - channel_dev.push(&ma, sizeof(mem_access_t)); - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - atomicAdd((unsigned long long*)&reported_dynamic_instr_counter, 1); - } -} -NVBIT_EXPORT_FUNC(instrument_mem); - - -extern "C" __device__ __noinline__ void instrument_inst(int pred, int opcode_id, - uint32_t vpc, int desReg, int srcReg1, int srcReg2, int srcReg3, int srcReg4, int srcNum) { - if (!pred) { - return; - } - - int active_mask = __ballot(1); - const int laneid = get_laneid(); - const int first_laneid = __ffs(active_mask) - 1; - - if (dynamic_instr_limit && total_dynamic_instr_counter >= dynamic_instr_limit) - if (first_laneid == laneid) { - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - return; - } - - - mem_access_t ma; - - int4 cta = get_ctaid(); - int uniqe_threadId = threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x; - ma.warpid_tb = uniqe_threadId/32; - - ma.cta_id_x = cta.x; - ma.cta_id_y = cta.y; - ma.cta_id_z = cta.z; - ma.warpid_sm = get_warpid(); - ma.opcode_id = opcode_id; - ma.is_mem = false; - ma.vpc = vpc; - - ma.GPRDst = desReg; - ma.numSrcs = srcNum; //this is the total src number including the register and others - ma.GPRSrcs[0] = srcReg1; - ma.GPRSrcs[1] = srcReg2; - ma.GPRSrcs[2] = srcReg3; - ma.GPRSrcs[3] = srcReg4; - - ma.active_mask = active_mask; - ma.sm_id = get_smid(); - - /* first active lane pushes information on the channel */ - if (first_laneid == laneid) { - channel_dev.push(&ma, sizeof(mem_access_t)); - atomicAdd((unsigned long long*)&total_dynamic_instr_counter, 1); - atomicAdd((unsigned long long*)&reported_dynamic_instr_counter, 1); - } -} - -NVBIT_EXPORT_FUNC(instrument_inst); - -void nvbit_at_init() { - setenv("CUDA_MANAGED_FORCE_DEVICE_ALLOC", "1", 1); - GET_VAR_INT( - instr_begin_interval, "INSTR_BEGIN", 0, - "Beginning of the instruction interval where to apply instrumentation"); - GET_VAR_INT( - instr_end_interval, "INSTR_END", UINT32_MAX, - "End of the instruction interval where to apply instrumentation"); - GET_VAR_LONG( - dynamic_instr_limit_input, "DYNAMIC_INSTR_LIMIT", 0, - "Limit of the number instructions to be printed, 0 means no limit"); - GET_VAR_INT(verbose, "TOOL_VERBOSE", 0, "Enable verbosity inside the tool"); - std::string pad(100, '-'); - printf("%s\n", pad.c_str()); -} - -/* instrument each memory instruction adding a call to the above instrumentation - * function */ -void nvbit_at_function_first_load(CUcontext ctx, CUfunction f) { - - dynamic_instr_limit = dynamic_instr_limit_input; - - const std::vector &instrs = nvbit_get_instrs(ctx, f); - if (verbose) { - printf("Inspecting function %s at address 0x%lx\n", - nvbit_get_func_name(ctx, f), nvbit_get_func_addr(f)); - } - - uint32_t cnt = 0; - /* iterate on all the static instructions in the function */ - for (auto instr : instrs) { - if (cnt < instr_begin_interval || cnt >= instr_end_interval ) { - cnt++; - continue; - } - //if (verbose) { - instr->printDecoded(); - //} - - if (opcode_to_id_map.find(instr->getOpcode()) == - opcode_to_id_map.end()) { - int opcode_id = opcode_to_id_map.size(); - opcode_to_id_map[instr->getOpcode()] = opcode_id; - id_to_opcode_map[opcode_id] = instr->getOpcode(); - } - - int opcode_id = opcode_to_id_map[instr->getOpcode()]; - - //TO DO: handle generic and TEX memory space - if(instr->isLoad() && !instr->isStore() && instr->getMemOpType() != Instr::CONSTANT) { //Mem load inst //ignore constant for now - assert(instr->getNumOperands() == 2); - - /* get the operand */ - const Instr::operand_t *dst = instr->getOperand(0); - const Instr::operand_t *src = instr->getOperand(1); - - assert(dst->type == Instr::REG); - assert(src->type == Instr::MREF); - - /* insert call to the instrumentation function with its - * arguments */ - nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - if (instr->isExtended()) { - nvbit_add_call_arg_reg_val(instr, (int)src->value[0] + 1); - } else { - nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ); - } - nvbit_add_call_arg_reg_val(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[1]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, -1); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)instr->getSize()); - } - else if(instr->isStore() && !instr->isLoad() && instr->getMemOpType() != Instr::CONSTANT) { //Mem store inst //ignore constant for now - assert(instr->getNumOperands() == 2); - - /* get the operand */ - const Instr::operand_t *dst = instr->getOperand(0); - const Instr::operand_t *src = instr->getOperand(1); - - assert(dst->type == Instr::MREF); - assert(src->type == Instr::REG); - - /* insert call to the instrumentation function with its - * arguments */ - nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - if (instr->isExtended()) { - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0] + 1); - } else { - nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ); - } - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[1]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, -1); - nvbit_add_call_arg_const_val32(instr, (int)instr->getSize()); - } - else if(instr->isLoad() && instr->isStore() && instr->getMemOpType() != Instr::CONSTANT) { //if it is load and store i.e. atomic inst - assert(instr->getNumOperands() == 2); - - /* get the operand */ - const Instr::operand_t *dst = instr->getOperand(0); - const Instr::operand_t *src = instr->getOperand(1); - - assert(dst->type == Instr::MREF); - assert(src->type == Instr::REG); - - /* insert call to the instrumentation function with its - * arguments */ - nvbit_insert_call(instr, "instrument_mem", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - if (instr->isExtended()) { - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0] + 1); - } else { - nvbit_add_call_arg_reg_val(instr, (int)Instr::RZ); - } - nvbit_add_call_arg_reg_val(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[1]); - nvbit_add_call_arg_const_val32(instr, (int)dst->value[0]); - nvbit_add_call_arg_const_val32(instr, (int)src->value[0]); - nvbit_add_call_arg_const_val32(instr, -1); - nvbit_add_call_arg_const_val32(instr, (int)instr->getSize()); - } - else //Other ALU, FP, DP insts - { - - nvbit_insert_call(instr, "instrument_inst", IPOINT_BEFORE); - nvbit_add_call_arg_pred_val(instr); - nvbit_add_call_arg_const_val32(instr, opcode_id); - nvbit_add_call_arg_const_val32(instr, (int)instr->getOffset()); - int srcNum = 0; - for (int i = 0; i < MAX_SRC+1; i++) { - /* get the operand "i" */ - if(i < instr->getNumOperands()) { - const Instr::operand_t *op = instr->getOperand(i); - if (op->type == Instr::REG) - nvbit_add_call_arg_const_val32(instr, (int)op->value[0]); - else - nvbit_add_call_arg_const_val32(instr, -1); - - srcNum++; - } - else - nvbit_add_call_arg_const_val32(instr, -1); - } - nvbit_add_call_arg_const_val32(instr, srcNum); - } - cnt++; - } -} - -__global__ void flush_channel() { - /* push memory access with negative cta id to communicate the kernel is - * completed */ - mem_access_t ma; - ma.cta_id_x = -1; - channel_dev.push(&ma, sizeof(mem_access_t)); - - /* flush channel */ - channel_dev.flush(); -} - -static FILE *resultsFile = NULL; -static FILE *kernelsFile= NULL; -static FILE *statsFile= NULL; -static int kernelid = 1; - -unsigned old_total_insts = 0; -unsigned old_total_reported_insts = 0; - - -void nvbit_at_cuda_event(CUcontext ctx, int is_exit, nvbit_api_cuda_t cbid, - const char *name, void *params, CUresult *pStatus) { - if (skip_flag) return; - - if (cbid == API_CUDA_cuLaunchKernel_ptsz || - cbid == API_CUDA_cuLaunchKernel) { - cuLaunchKernel_params *p = (cuLaunchKernel_params *)params; - - if (!is_exit) { - - - if (mkdir("traces", S_IRWXU | S_IRWXG | S_IROTH | S_IXOTH) == -1) { - if( errno == EEXIST ) { - // alredy exists - } else { - // something else - std::cout << "cannot create folder error:" << strerror(errno) << std::endl; - return; - } - } - - int nregs; - _cuda_safe( - cuFuncGetAttribute(&nregs, CU_FUNC_ATTRIBUTE_NUM_REGS, p->f)); - - int shmem_static_nbytes; - _cuda_safe(cuFuncGetAttribute(&shmem_static_nbytes, - CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, - p->f)); - - - - std::string func_name(nvbit_get_func_name(ctx, p->f)); - std::string::size_type end_pos = func_name.find('('); - if (end_pos != std::string::npos) - { - // std::string::size_type pos = func_name.find('<'); - //if (pos != std::string::npos) - // end_pos = pos; - - //std::string::size_type start_pos = func_name.find(' '); - //if (start_pos == std::string::npos) - // start_pos = 0; - //else - // start_pos++; - - func_name = func_name.substr(0, end_pos); - } - - char buffer[1024]; - sprintf (buffer, "./traces/%d-%s.trace", kernelid, func_name.c_str()); - - resultsFile = fopen(buffer, "w"); - - printf("Writing results to %s\n", buffer); - - fprintf(resultsFile, "-kernel name = %s", nvbit_get_func_name(ctx, p->f)); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-kernel id = %d", kernelid); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-grid dim = (%d,%d,%d)", p->gridDimX, p->gridDimY, p->gridDimZ); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-block dim = (%d,%d,%d)", p->blockDimX, p->blockDimY, p->blockDimZ); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-shmem = %d", shmem_static_nbytes + p->sharedMemBytes); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-nregs = %d", nregs); - fprintf(resultsFile, "\n"); - fprintf(resultsFile, "-cuda stream id = %d", (uint64_t)p->hStream); - fprintf(resultsFile, "\n\n"); - - fprintf(resultsFile, "#traces format = threadblock_x threadblock_y threadblock_z warpid_tb sm_id warpid_sm PC mask dest_num reg_dests opcode src_num reg_srcs mem_width mem_addresses"); - fprintf(resultsFile, "\n"); - - if (kernelid == 1) { - kernelsFile = fopen("./traces/kernelslist", "w"); - statsFile = fopen("./traces/stats.csv", "w"); - fprintf(statsFile, "kernel name,total_insts,total_reported_insts\n"); - } - else { - kernelsFile = fopen("./traces/kernelslist", "a"); - statsFile = fopen("./traces/stats.csv", "a"); - } - - sprintf (buffer, "%d-%s.trace", kernelid, func_name.c_str()); - fprintf(kernelsFile, buffer); - fprintf(kernelsFile, "\n"); - fclose(kernelsFile); - - fprintf(statsFile, buffer); - fprintf(statsFile, ","); - - kernelid++; - recv_thread_receiving = true; - - } else { - /* make sure current kernel is completed */ - cudaDeviceSynchronize(); - assert(cudaGetLastError() == cudaSuccess); - - /* make sure we prevent re-entry on the nvbit_callback when issuing - * the flush_channel kernel */ - skip_flag = true; - - /* issue flush of channel so we are sure all the memory accesses - * have been pushed */ - flush_channel<<<1, 1>>>(); - cudaDeviceSynchronize(); - assert(cudaGetLastError() == cudaSuccess); - - /* unset the skip flag */ - skip_flag = false; - - /* wait here until the receiving thread has not finished with the - * current kernel */ - while (recv_thread_receiving) { - pthread_yield(); - } - - unsigned total_insts_per_kernel = total_dynamic_instr_counter - old_total_insts; - old_total_insts = total_dynamic_instr_counter; - - unsigned reported_insts_per_kernel = reported_dynamic_instr_counter - old_total_reported_insts; - old_total_reported_insts = reported_dynamic_instr_counter; - - fprintf(statsFile, ""); - fprintf(statsFile, "%d,%d",total_insts_per_kernel,reported_insts_per_kernel); - fprintf(statsFile, "\n"); - - - fclose(resultsFile); - fclose(statsFile); - } - } -} - -bool is_number(const std::string& s) -{ - std::string::const_iterator it = s.begin(); - while (it != s.end() && std::isdigit(*it)) ++it; - return !s.empty() && it == s.end(); -} - -void *recv_thread_fun(void *) { - char *recv_buffer = (char *)malloc(CHANNEL_SIZE); - - while (recv_thread_started) { - uint32_t num_recv_bytes = 0; - if (recv_thread_receiving && - (num_recv_bytes = channel_host.recv(recv_buffer, CHANNEL_SIZE)) > - 0) { - uint32_t num_processed_bytes = 0; - while (num_processed_bytes < num_recv_bytes) { - mem_access_t *ma = - (mem_access_t *)&recv_buffer[num_processed_bytes]; - - /* when we get this cta_id_x it means the kernel has completed - */ - if (ma->cta_id_x == -1) { - recv_thread_receiving = false; - break; - } - - fprintf(resultsFile, "%d ", ma->cta_id_x); - fprintf(resultsFile, "%d ", ma->cta_id_y); - fprintf(resultsFile, "%d ", ma->cta_id_z); - fprintf(resultsFile, "%d ", ma->warpid_tb); - fprintf(resultsFile, "%d ", ma->sm_id); - fprintf(resultsFile, "%d ", ma->warpid_sm); - fprintf(resultsFile, "0x%016lx ", ma->vpc); // Print the virtual PC. - fprintf(resultsFile, "%-8.8" PRIx32 " ", ma->active_mask); - if(ma->GPRDst >= 0) { - fprintf(resultsFile, "1 "); - fprintf(resultsFile, "R%d ", ma->GPRDst); - } - else - fprintf(resultsFile, "0 "); - - // Print the opcode. - fprintf(resultsFile, "%s ", id_to_opcode_map[ma->opcode_id].c_str()); - unsigned src_count=0; - for (int s = 0; s < MAX_SRC; s++) // GPR srcs count. - if(ma->GPRSrcs[s] >= 0) src_count++; - fprintf(resultsFile, "%d ", src_count); - - for (int s = 0; s < MAX_SRC; s++) // GPR srcs. - if(ma->GPRSrcs[s] >= 0) fprintf(resultsFile, "R%d ", ma->GPRSrcs[s]); - - //print addresses - std::bitset<32> mask(ma->active_mask); - if(ma->is_mem) { - //fprintf(resultsFile, "%d ", ma->width); - std::istringstream iss(id_to_opcode_map[ma->opcode_id]); - std::vector tokens; - std::string token; - while (std::getline(iss, token, '.')) { - if (!token.empty()) - tokens.push_back(token); - } - if (tokens.size()>=3){ - if (is_number(tokens[2])){ - fprintf(resultsFile, "%d ", (std::stoi(tokens[2],nullptr)/8)); - } - else{ - fprintf(resultsFile, "%d ", 4); - } - } - else{ - fprintf(resultsFile, "%d ", 4); - } - - for (int s = 0; s < 32; s++) - if(mask.test(s)) - fprintf(resultsFile, "0x%016lx ", ma->addrs[s]); - } - else - { - fprintf(resultsFile, "0 "); - } - - fprintf(resultsFile, "\n"); - - num_processed_bytes += sizeof(mem_access_t); - } - } - } - free(recv_buffer); - return NULL; -} - -void nvbit_at_ctx_init(CUcontext ctx) { - recv_thread_started = true; - channel_host.init(0, CHANNEL_SIZE, &channel_dev, NULL); - pthread_create(&recv_thread, NULL, recv_thread_fun, NULL); -} - -void nvbit_at_ctx_term(CUcontext ctx) { - if (recv_thread_started) { - recv_thread_started = false; - pthread_join(recv_thread, NULL); - } -} diff --git a/src/trace-driven/traces-generator/README b/src/trace-driven/traces-generator/README deleted file mode 100644 index d79d075..0000000 --- a/src/trace-driven/traces-generator/README +++ /dev/null @@ -1 +0,0 @@ -TO DO diff --git a/src/trace-driven/traces-generator/traces-post-processing/Makefile b/src/trace-driven/traces-generator/traces-post-processing/Makefile deleted file mode 100755 index 63c3f38..0000000 --- a/src/trace-driven/traces-generator/traces-post-processing/Makefile +++ /dev/null @@ -1,11 +0,0 @@ -TARGET := post-traces-processing - -$(TARGET): post-traces-processing.cpp - g++ -o $@ $^ - -run: $(TARGET) - ./$(TARGET) - -clean: - rm -f $(TARGET) *.o - diff --git a/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp b/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp deleted file mode 100644 index 0dbb4e1..0000000 --- a/src/trace-driven/traces-generator/traces-post-processing/post-traces-processing.cpp +++ /dev/null @@ -1,195 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include -using namespace std; - -struct threadblock_info -{ - bool initialized; - unsigned tb_id_x, tb_id_y, tb_id_z; - vector< vector< string > > warp_insts_array; - threadblock_info() { - initialized = false; - tb_id_x = tb_id_y = tb_id_z = 0; - } -}; - -void group_per_block(const char* filepath); -void group_per_core(const char* filepath); - -////////////////////////////////////////////////////////////////////////////////////////////////////////////////// -int main(int argc, char** argv) -{ - - string kernellist_filepath; - bool is_per_core; - if(argc == 1) - { - cout << "File path is missing\n"; - return 0; - } else if(argc == 2) - { - kernellist_filepath = argv[1]; - is_per_core = true; - - } else if(argc == 3) { - kernellist_filepath = argv[1]; - is_per_core = bool(argv[2]); - } - else { - cout << "Too Many Arguemnts!\n"; - return 0; - } - - ifstream ifs; - ofstream ofs; - - ifs.open(kernellist_filepath.c_str()); - ofs.open((string(kernellist_filepath) + ".g").c_str()); - - if (!ifs.is_open()) { - cout << "Unable to open file: " < insts; - unsigned grid_dim_x, grid_dim_y, grid_dim_z, tb_dim_x, tb_dim_y, tb_dim_z; - unsigned tb_id_x, tb_id_y, tb_id_z, tb_id, warpid_tb; - string line; - stringstream ss; - string string1, string2; - bool found_grid_dim = false, found_block_dim = false; - - while(!ifs.eof()) { - getline(ifs, line); - - if (line.length() == 0 || line[0] == '#') { - ofs<>string1>>string2; - if(string1 == "grid" && string2 == "dim") { - sscanf(line.c_str(), "-grid dim = (%d,%d,%d)", &grid_dim_x, &grid_dim_y, &grid_dim_z); - found_grid_dim = true; - } - else if (string1 == "block" && string2 == "dim") { - sscanf(line.c_str(), "-block dim = (%d,%d,%d)", &tb_dim_x, &tb_dim_y, &tb_dim_z); - found_block_dim = true; - } - - if(found_grid_dim && found_block_dim) { - insts.resize(grid_dim_x*grid_dim_y*grid_dim_z); - for(unsigned i = 0; i>tb_id_x>>tb_id_y>>tb_id_z>>warpid_tb; - tb_id = tb_id_z * grid_dim_y * grid_dim_x + tb_id_y * grid_dim_x + tb_id_x; - if(!insts[tb_id].initialized) { - insts[tb_id].tb_id_x = tb_id_x; - insts[tb_id].tb_id_y = tb_id_y; - insts[tb_id].tb_id_z = tb_id_z; - insts[tb_id].initialized = true; - } - insts[tb_id].warp_insts_array[warpid_tb].push_back(line); - } - - } - - - for(unsigned i=0; i 0) { - ofs< -#include - -//TO DO: moving this to a yml or def files - - -#define TURING_BINART_VERSION 72 - -///Tuing SM_72 ISA -//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html -static const std::unordered_map Turing_OpcodeMap = { - -//TO fill - -}; - -#endif diff --git a/src/trace-driven/volta_opcode.h b/src/trace-driven/volta_opcode.h deleted file mode 100644 index 3d03201..0000000 --- a/src/trace-driven/volta_opcode.h +++ /dev/null @@ -1,176 +0,0 @@ -//developed by Mahmoud Khairy, Purdue Univ -//abdallm@purdue.edu - -#ifndef VOLTA_OPCODE_H -#define VOLTA_OPCODE_H - -#include "../abstract_hardware_model.h" -#include "trace_opcode.h" -#include -#include - -#define VOLTA_BINART_VERSION 70 -#define VOLTA_SHARED_MEMORY_VIRTIAL_ADDRESS_START 0x00007f2c60000000 - -//TO DO: moving this to a yml or def files - -///Volta SM_70 ISA -//see: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html -static const std::unordered_map Volta_OpcodeMap = { - //Floating Point 32 Instructions - {"FADD", OpcodeChar(OP_FADD, SP_OP)}, - {"FADD32I", OpcodeChar(OP_FADD32I, SP_OP)}, - {"FCHK", OpcodeChar(OP_FCHK, SP_OP)}, - {"FFMA32I", OpcodeChar(OP_FFMA32I, SP_OP)}, - {"FFMA", OpcodeChar(OP_FFMA, SP_OP)}, - {"FMNMX", OpcodeChar(OP_FMNMX, SP_OP)}, - {"FMUL", OpcodeChar(OP_FMUL, SP_OP)}, - {"FMUL32I", OpcodeChar(OP_FMUL32I, SP_OP)}, - {"FSEL", OpcodeChar(OP_FSEL, SP_OP)}, - {"FSET", OpcodeChar(OP_FSET, SP_OP)}, - {"FSETP", OpcodeChar(OP_FSETP, SP_OP)}, - {"FSWZADD", OpcodeChar(OP_FSWZADD, SP_OP)}, - //SFU - {"MUFU", OpcodeChar(OP_MUFU, SFU_OP)}, - - //Floating Point 16 Instructions - {"HADD2", OpcodeChar(OP_HADD2, SP_OP)}, - {"HADD2_32I", OpcodeChar(OP_HADD2_32I, SP_OP)}, - {"HFMA2", OpcodeChar(OP_HFMA2, SP_OP)}, - {"HFMA2_32I", OpcodeChar(OP_HFMA2_32I, SP_OP)}, - {"HMUL2", OpcodeChar(OP_HMUL2, SP_OP)}, - {"HMUL2_32I", OpcodeChar(OP_HMUL2_32I, SP_OP)}, - {"HSET2", OpcodeChar(OP_HSET2, SP_OP)}, - {"HSETP2", OpcodeChar(OP_HSETP2, SP_OP)}, - - //Tensor Core Instructions - {"HMMA", OpcodeChar(OP_HMMA, TENSOR_CORE_OP)}, - - //Double Point Instructions - {"DADD", OpcodeChar(OP_DADD, DP_OP)}, - {"DFMA", OpcodeChar(OP_DFMA, DP_OP)}, - {"DMUL", OpcodeChar(OP_DMUL, DP_OP)}, - {"DSETP", OpcodeChar(OP_DSETP, DP_OP)}, - - //Integer Instructions - {"BMSK", OpcodeChar(OP_BMSK, INTP_OP)}, - {"BREV", OpcodeChar(OP_BREV, INTP_OP)}, - {"FLO", OpcodeChar(OP_FLO, INTP_OP)}, - {"IABS", OpcodeChar(OP_IABS, INTP_OP)}, - {"IADD", OpcodeChar(OP_IADD, INTP_OP)}, - {"IADD3", OpcodeChar(OP_IADD3, INTP_OP)}, - {"IADD32I", OpcodeChar(OP_IADD32I, INTP_OP)}, - {"IDP", OpcodeChar(OP_IDP, INTP_OP)}, - {"IDP4A", OpcodeChar(OP_IDP4A, INTP_OP)}, - {"IMAD", OpcodeChar(OP_IMAD, INTP_OP)}, - {"IMMA", OpcodeChar(OP_IMMA, INTP_OP)}, - {"IMNMX", OpcodeChar(OP_IMNMX, INTP_OP)}, - {"IMUL", OpcodeChar(OP_IMUL, INTP_OP)}, - {"IMUL32I", OpcodeChar(OP_IMUL32I, INTP_OP)}, - {"ISCADD", OpcodeChar(OP_ISCADD, INTP_OP)}, - {"ISCADD32I", OpcodeChar(OP_ISCADD32I, INTP_OP)}, - {"ISETP", OpcodeChar(OP_ISETP, INTP_OP)}, - {"LEA", OpcodeChar(OP_LEA, INTP_OP)}, - {"LOP", OpcodeChar(OP_LOP, INTP_OP)}, - {"LOP3", OpcodeChar(OP_LOP3, INTP_OP)}, - {"LOP32I", OpcodeChar(OP_LOP32I, INTP_OP)}, - {"POPC", OpcodeChar(OP_POPC, INTP_OP)}, - {"SHF", OpcodeChar(OP_SHF, INTP_OP)}, - {"SHR", OpcodeChar(OP_SHR, INTP_OP)}, - {"VABSDIFF", OpcodeChar(OP_VABSDIFF, INTP_OP)}, - {"VABSDIFF4", OpcodeChar(OP_VABSDIFF4, INTP_OP)}, - - //Conversion Instructions - {"F2F", OpcodeChar(OP_F2F, ALU_OP)}, - {"F2I", OpcodeChar(OP_F2I, ALU_OP)}, - {"I2F", OpcodeChar(OP_I2F, ALU_OP)}, - {"I2I", OpcodeChar(OP_I2I, ALU_OP)}, - {"I2IP", OpcodeChar(OP_I2IP, ALU_OP)}, - {"FRND", OpcodeChar(OP_FRND, ALU_OP)}, - - //Movement Instructions - {"MOV", OpcodeChar(OP_MOV, ALU_OP)}, - {"MOV32I", OpcodeChar(OP_MOV32I, ALU_OP)}, - {"PRMT", OpcodeChar(OP_PRMT, ALU_OP)}, - {"SEL", OpcodeChar(OP_SEL, ALU_OP)}, - {"SGXT", OpcodeChar(OP_SGXT, ALU_OP)}, - {"SHFL", OpcodeChar(OP_SHFL, ALU_OP)}, - - //Predicate Instructions - {"PLOP3", OpcodeChar(OP_PLOP3, ALU_OP)}, - {"PSETP", OpcodeChar(OP_PSETP, ALU_OP)}, - {"P2R", OpcodeChar(OP_P2R, ALU_OP)}, - {"R2P", OpcodeChar(OP_R2P, ALU_OP)}, - - //Load/Store Instructions - {"LD", OpcodeChar(OP_LD, LOAD_OP)}, - //For now, we ignore constant loads, consider it as ALU_OP, TO DO - {"LDC", OpcodeChar(OP_LDC, ALU_OP)}, - {"LDG", OpcodeChar(OP_LDG, LOAD_OP)}, - {"LDL", OpcodeChar(OP_LDL, LOAD_OP)}, - {"LDS", OpcodeChar(OP_LDS, LOAD_OP)}, - {"ST", OpcodeChar(OP_ST, STORE_OP)}, - {"STG", OpcodeChar(OP_STG, STORE_OP)}, - {"STL", OpcodeChar(OP_STL, STORE_OP)}, - {"STS", OpcodeChar(OP_STS, STORE_OP)}, - {"MATCH", OpcodeChar(OP_MATCH, ALU_OP)}, - {"QSPC", OpcodeChar(OP_QSPC, ALU_OP)}, - {"ATOM", OpcodeChar(OP_ATOM, STORE_OP)}, - {"ATOMS", OpcodeChar(OP_ATOMS, STORE_OP)}, - {"ATOMG", OpcodeChar(OP_ATOMG, STORE_OP)}, - {"RED", OpcodeChar(OP_RED, STORE_OP)}, - {"CCTL", OpcodeChar(OP_CCTL, ALU_OP)}, - {"CCTLL", OpcodeChar(OP_CCTLL, ALU_OP)}, - {"ERRBAR", OpcodeChar(OP_ERRBAR, ALU_OP)}, - {"MEMBAR", OpcodeChar(OP_MEMBAR, MEMORY_BARRIER_OP)}, - {"CCTLT", OpcodeChar(OP_CCTLT, ALU_OP)}, - - //Texture Instructions - //For now, we ignore texture loads, consider it as ALU_OP - {"TEX", OpcodeChar(OP_TEX, ALU_OP)}, - {"TLD", OpcodeChar(OP_TLD, ALU_OP)}, - {"TLD4", OpcodeChar(OP_TLD4, ALU_OP)}, - {"TMML", OpcodeChar(OP_TMML, ALU_OP)}, - {"TXD", OpcodeChar(OP_TXD, ALU_OP)}, - {"TXQ", OpcodeChar(OP_TXQ, ALU_OP)}, - - //Control Instructions - {"BMOV", OpcodeChar(OP_BMOV, BRANCH_OP)}, - {"BPT", OpcodeChar(OP_BPT, BRANCH_OP)}, - {"BRA", OpcodeChar(OP_BRA, BRANCH_OP)}, - {"BREAK", OpcodeChar(OP_BREAK, BRANCH_OP)}, - {"BRX", OpcodeChar(OP_BRX, BRANCH_OP)}, - {"BSSY", OpcodeChar(OP_BSSY, BRANCH_OP)}, - {"BSYNC", OpcodeChar(OP_BSYNC, BRANCH_OP)}, - {"CALL", OpcodeChar(OP_CALL, CALL_OPS)}, - {"EXIT", OpcodeChar(OP_EXIT, EXIT_OPS)}, - {"JMP", OpcodeChar(OP_JMP, BRANCH_OP)}, - {"JMX", OpcodeChar(OP_JMX, BRANCH_OP)}, - {"KILL", OpcodeChar(OP_KILL, BRANCH_OP)}, - {"NANOSLEEP", OpcodeChar(OP_NANOSLEEP, BRANCH_OP)}, - {"RET", OpcodeChar(OP_RET, RET_OPS)}, - {"RPCMOV", OpcodeChar(OP_RPCMOV, BRANCH_OP)}, - {"RTT", OpcodeChar(OP_RTT, RET_OPS)}, - {"WARPSYNC", OpcodeChar(OP_WARPSYNC, BRANCH_OP)}, - {"YIELD", OpcodeChar(OP_YIELD, BRANCH_OP)}, - - //Miscellaneous Instructions - {"B2R", OpcodeChar(OP_B2R, ALU_OP)}, - {"BAR", OpcodeChar(OP_BAR, BARRIER_OP)}, - {"CS2R", OpcodeChar(OP_CS2R, ALU_OP)}, - {"CSMTEST", OpcodeChar(OP_CSMTEST, ALU_OP)}, - {"DEPBAR", OpcodeChar(OP_DEPBAR, ALU_OP)}, - {"GETLMEMBASE", OpcodeChar(OP_GETLMEMBASE, ALU_OP)}, - {"LEPC", OpcodeChar(OP_LEPC ,ALU_OP)}, - {"NOP", OpcodeChar(OP_NOP ,ALU_OP)}, - {"PMTRIG", OpcodeChar(OP_PMTRIG, ALU_OP)}, - {"R2B", OpcodeChar(OP_R2B, ALU_OP)}, - {"S2R", OpcodeChar(OP_S2R, ALU_OP)}, - {"SETCTAID", OpcodeChar(OP_SETCTAID, ALU_OP)}, - {"SETLMEMBASE", OpcodeChar(OP_SETLMEMBASE, ALU_OP)}, - {"VOTE", OpcodeChar(OP_VOTE, ALU_OP)}, - {"VOTE_VTG", OpcodeChar(OP_VOTE_VTG, ALU_OP)}, - -}; - -#endif -- cgit v1.3