summaryrefslogtreecommitdiff
path: root/cuobjdump_to_ptxplus
diff options
context:
space:
mode:
Diffstat (limited to 'cuobjdump_to_ptxplus')
-rw-r--r--cuobjdump_to_ptxplus/Makefile190
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInst.cc193
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInst.h43
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInstList.cc452
-rw-r--r--cuobjdump_to_ptxplus/cuobjdumpInstList.h108
-rw-r--r--cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc4
-rw-r--r--cuobjdump_to_ptxplus/elf.l2
-rw-r--r--cuobjdump_to_ptxplus/header.l124
-rw-r--r--cuobjdump_to_ptxplus/header.y116
-rw-r--r--cuobjdump_to_ptxplus/ptx_parser.h25
-rw-r--r--cuobjdump_to_ptxplus/sass.l2
-rw-r--r--cuobjdump_to_ptxplus/stringList.cc128
-rw-r--r--cuobjdump_to_ptxplus/stringList.h64
13 files changed, 1191 insertions, 260 deletions
diff --git a/cuobjdump_to_ptxplus/Makefile b/cuobjdump_to_ptxplus/Makefile
index 2f3c202..7f58f69 100644
--- a/cuobjdump_to_ptxplus/Makefile
+++ b/cuobjdump_to_ptxplus/Makefile
@@ -1,11 +1,9 @@
-BOOST_ROOT ?= /usr/include/boost
-
LEX = flex
LEXFLAGS = -B
YACC = bison
YFLAGS = -t -d -v --report=all
CC = gcc
-CCFLAGS = -ggdb -I${BOOST_ROOT} -lboost_regex -lboost_filesystem -fPIC
+CCFLAGS = -ggdb -fPIC
CXX = g++
CXXFLAGS = ${CCFLAGS}
LD = g++
@@ -14,71 +12,18 @@ LDFLAGS = ${CCFLAGS}
SASS_PARSER_OBJECTS = sass_lexer.o sass_parser.o
ELF_PARSER_OBJECTS = elf_lexer.o elf_parser.o
-HEADER_PARSER_OBJECTS = headerParser.o headerLexer.o
+HEADER_PARSER_OBJECTS = header_parser.o header_lexer.o
PTX_PARSER_OBJECTS = ptx.tab.o lex.ptx_.o
all: cuobjdump_to_ptxplus
-cuobjdump_to_ptxplus: stringList.o decudaInst.o decudaInstList.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o $(PTX_PARSER_OBJECTS) $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) $(HEADER_PARSER_OBJECTS)
- ${LD} ${LDFLAGS} -ocuobjdump_to_ptxplus stringList.o decudaInst.o decudaInstList.o headerParser.o headerLexer.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o lex.ptx_.o ptx.tab.o $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS)
-
-stringList.o: ../decuda_to_ptxplus/stringList.cc ../decuda_to_ptxplus/stringList.h
- ${CXX} ${CXXFLAGS} -c -ostringList.o ../decuda_to_ptxplus/stringList.cc
-
-decudaInst.o: ../decuda_to_ptxplus/decudaInst.cc ../decuda_to_ptxplus/decudaInst.h ../decuda_to_ptxplus/stringList.h
- ${CXX} ${CXXFLAGS} -c -odecudaInst.o ../decuda_to_ptxplus/decudaInst.cc
-
-decudaInstList.o: ../decuda_to_ptxplus/decudaInstList.cc ../decuda_to_ptxplus/decudaInstList.h ../decuda_to_ptxplus/decudaInst.h
- ${CXX} ${CXXFLAGS} -c -odecudaInstList.o ../decuda_to_ptxplus/decudaInstList.cc
-
-cuobjdumpInst.o: cuobjdumpInst.cc cuobjdumpInst.h ../decuda_to_ptxplus/decudaInst.h ./decudaInstList.h
- ${CXX} ${CXXFLAGS} -c -ocuobjdumpInst.o cuobjdumpInst.cc
-
-cuobjdumpInstList.o: cuobjdumpInstList.cc cuobjdumpInstList.h ../decuda_to_ptxplus/decudaInstList.h
- ${CXX} ${CXXFLAGS} -c -ocuobjdumpInstList.o cuobjdumpInstList.cc
-
-#sassParser.cpp: sassParser.y ../distribution/decuda_to_ptxplus/decudaInstList.h cuobjdumpInstList.h
-# ${YACC} ${YFLAGS} -osassParser.cpp sassParser.y
-
-#sassParser.hpp: sassParser.cpp
-
-#sassLexer.cpp: sassLexer.l sassParser.hpp
-# ${LEX} ${LEXFLAGS} -osassLexer.cpp sassLexer.l
-
-headerParser.cpp : ../decuda_to_ptxplus/header.y ../decuda_to_ptxplus/decudaInstList.h
- ${YACC} ${YFLAGS} --name-prefix=header_ -oheaderParser.cpp ../decuda_to_ptxplus/header.y
-
-headerParser.hpp : headerParser.cpp
-
-headerLexer.cpp : ../decuda_to_ptxplus/header.l headerParser.hpp
- ${LEX} ${LEXFLAGS} -oheaderLexer.cpp ../decuda_to_ptxplus/header.l
-
-#sassLexer.o: sassLexer.cpp
-# ${CXX} ${CXXFLAGS} -c -Igenerated -osassLexer.o sassLexer.cpp
-#
-#sassParser.o: sassParser.cpp ../distribution/decuda_to_ptxplus/stringList.h ../distribution/decuda_to_ptxplus/decudaInst.h ../distribution/decuda_to_ptxplus/decudaInstList.h cuobjdumpInstList.h
-# ${CXX} ${CXXFLAGS} -c -Igenerated -osassParser.o sassParser.cpp
-
-headerLexer.o : headerLexer.cpp
- ${CXX} ${CXXFLAGS} -c -Igenerated -oheaderLexer.o headerLexer.cpp
+MAKEFLAGS += --no-builtin-rules
-headerParser.o : headerParser.cpp ../decuda_to_ptxplus/stringList.h ../decuda_to_ptxplus/decudaInst.h ./decudaInstList.h
- ${CXX} ${CXXFLAGS} -c -Igenerated -oheaderParser.o headerParser.cpp
+.SUFFIXES:
+.SECONDARY:
-decudaInstList.h : ../decuda_to_ptxplus/decudaInstList.h ./decudaInst.h
- rm -f decudaInstList.h
- ln -s ../decuda_to_ptxplus/decudaInstList.h decudaInstList.h
-
-decudaInst.h : ../decuda_to_ptxplus/decudaInst.h ./stringList.h
- rm -f decudaInst.h
- ln -s ../decuda_to_ptxplus/decudaInst.h decudaInst.h
-
-stringList.h : ../decuda_to_ptxplus/stringList.h
- rm -f stringList.h
- ln -s ../decuda_to_ptxplus/stringList.h stringList.h
-
-cuobjdump_to_ptxplus.o: cuobjdump_to_ptxplus.cc ../decuda_to_ptxplus/decudaInstList.h cuobjdumpInstList.h
- ${CXX} ${CXXFLAGS} -c -ocuobjdump_to_ptxplus.o cuobjdump_to_ptxplus.cc
+cuobjdump_to_ptxplus: stringList.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o $(PTX_PARSER_OBJECTS) $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) $(HEADER_PARSER_OBJECTS)
+ ${LD} ${LDFLAGS} -o $@ stringList.o cuobjdumpInst.o cuobjdumpInstList.o cuobjdump_to_ptxplus.o $(PTX_PARSER_OBJECTS) $(SASS_PARSER_OBJECTS) $(ELF_PARSER_OBJECTS) $(HEADER_PARSER_OBJECTS)
lex.ptx_.c : ../src/cuda-sim/ptx.l
${LEX} ${LEXFLAGS} ../src/cuda-sim/ptx.l
@@ -97,37 +42,27 @@ lex.ptx_.o : lex.ptx_.c ptx.tab.h
ptx.tab.o : ptx.tab.c ptx_parser.h
${CXX} ${CXXFLAGS} -c ptx.tab.c
-.PRECIOUS: %.c %.h %_lexer.c %_parser.c %_parser.h %_parser.cpp %_parser.hpp %_lexer.cpp
-%_lexer.c: %.l %_parser.h
+%_lexer.cc: %.l %_parser.hh
$(LEX) $(LEXFLAGS) -P $*_ -o$@ $<
-%_lexer.cpp: %.l %_parser.hpp
- $(LEX) $(LEXFLAGS) -P $*_ -o$@ $<
-
-%_parser.c: %.y
- $(YACC) $(YFLAGS) -p $*_ -o$@ $<
-
-%_parser.cpp: %.y
+%_parser.cc: %.y
$(YACC) $(YFLAGS) -p $*_ -o$@ $<
# The ':' is needed otherwise make gets confused and doesn't understand that it knows how to make %Parser.hpp from %Parser.cpp.
# : in bash is equivalent to NOP (no operation)
-%_parser.h: %_parser.c
+%_parser.hh: %_parser.cc
:
-%_parser.hpp: %_parser.cpp
- :
+%.o: %.cc
+ $(CXX) ${CXXFLAGS} -c -o $@ $<
-cuobjdump_parser.o: cuobjdump_parser.c
- ${CC} ${CCFLAGS} -c -o $@ $<
+%.o: %.cpp %.h
+ $(CXX) ${CXXFLAGS} -c -o $@ $<
-cuobjdump_lexer.o: cuobjdump_lexer.c
+%.o: %.c %.h
${CC} ${CCFLAGS} -c -o $@ $<
-%.o: %.cpp
- $(CXX) ${CXXFLAGS} -c -o $@ $<
-
SRCS = $(shell ls *.cc)
depend:
@@ -136,22 +71,53 @@ depend:
clean:
rm -f ./opcodes.def ./opcodes.h
rm -f ./lex.ptx_.c ./ptx.tab.c ./ptx.tab.h ./lex.ptx_.o ./ptx.tab.o ptx.output
- rm -f ./decudaInstList.h ./decudaInst.h ./stringList.h
rm -f ./headerParser.o ./headerLexer.o ./headerParser.cpp ./headerParser.hpp ./headerLexer.cpp ./headerParser.output
- rm -f ./stringList.o ./decudaInst.o ./decudaInstList.o
+ rm -f ./header_parser.output ./header_parser.o header_parser.cc header_parser.hh header_lexer.cc header_lexer.hh
+ rm -f ./stringList.o
rm -f ./cuobjdumpInstList.o ./cuobjdumpInst.o
rm -f ./sass_lexer.o ./sass_parser.o ./cuobjdump_to_ptxplus.o
- rm -f ./sass_parser.hpp ./sass_parser.output
+ rm -f ./sass_parser.hpp ./sass_parser.output ./sass_parser.h
rm -f ./sass_parser.cpp ./sass_lexer.cpp
rm -f ./cuobjdump_to_ptxplus
rm -f ./cuobjdump_parser.o ./cuobjdump_parser.c ./cuobjdump_parser.output ./cuobjdump_parser.h
rm -f ./cuobjdump_lexer.o ./cuobjdump_lexer.c
rm -f ./elf_lexer.o ./elf_parser.o ./elf_lexer.cpp ./elf_parser.cpp ./elf_parser.hpp ./elf_parser.output
+ rm -f ./header_lexer.o header_lexer.c header_lexer.h sass_lexer.c sass_lexer.h sass_parser.c
+ rm -f ./sass_lexer.cc ./sass_parser.cc ./sass_parser.hh ./sass_lexer.hh
# DO NOT DELETE
-cuobjdumpInst.o: cuobjdumpInst.h
+cuobjdumpInst.o: cuobjdumpInst.h /usr/include/string.h
+cuobjdumpInst.o: /usr/include/features.h /usr/include/sys/cdefs.h
+cuobjdumpInst.o: /usr/include/bits/wordsize.h /usr/include/gnu/stubs.h
+cuobjdumpInst.o: /usr/include/gnu/stubs-64.h /usr/include/xlocale.h
+cuobjdumpInst.o: /usr/include/stdio.h /usr/include/bits/types.h
+cuobjdumpInst.o: /usr/include/bits/typesizes.h /usr/include/libio.h
+cuobjdumpInst.o: /usr/include/_G_config.h /usr/include/wchar.h
+cuobjdumpInst.o: /usr/include/bits/stdio_lim.h
+cuobjdumpInst.o: /usr/include/bits/sys_errlist.h /usr/include/stdlib.h
+cuobjdumpInst.o: /usr/include/sys/types.h /usr/include/time.h
+cuobjdumpInst.o: /usr/include/endian.h /usr/include/bits/endian.h
+cuobjdumpInst.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h
+cuobjdumpInst.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
+cuobjdumpInst.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
+cuobjdumpInst.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
+cuobjdumpInst.o: stringList.h
cuobjdumpInstList.o: cuobjdumpInstList.h cuobjdumpInst.h
-cuobjdump_to_ptxplus.o: cuobjdumpInstList.h cuobjdumpInst.h
+cuobjdumpInstList.o: /usr/include/string.h /usr/include/features.h
+cuobjdumpInstList.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+cuobjdumpInstList.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+cuobjdumpInstList.o: /usr/include/xlocale.h /usr/include/stdio.h
+cuobjdumpInstList.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+cuobjdumpInstList.o: /usr/include/libio.h /usr/include/_G_config.h
+cuobjdumpInstList.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+cuobjdumpInstList.o: /usr/include/bits/sys_errlist.h /usr/include/stdlib.h
+cuobjdumpInstList.o: /usr/include/sys/types.h /usr/include/time.h
+cuobjdumpInstList.o: /usr/include/endian.h /usr/include/bits/endian.h
+cuobjdumpInstList.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h
+cuobjdumpInstList.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
+cuobjdumpInstList.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
+cuobjdumpInstList.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
+cuobjdumpInstList.o: stringList.h
cuobjdump_to_ptxplus.o: /usr/include/stdio.h /usr/include/features.h
cuobjdump_to_ptxplus.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
cuobjdump_to_ptxplus.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
@@ -159,4 +125,56 @@ cuobjdump_to_ptxplus.o: /usr/include/bits/types.h
cuobjdump_to_ptxplus.o: /usr/include/bits/typesizes.h /usr/include/libio.h
cuobjdump_to_ptxplus.o: /usr/include/_G_config.h /usr/include/wchar.h
cuobjdump_to_ptxplus.o: /usr/include/bits/stdio_lim.h
-cuobjdump_to_ptxplus.o: /usr/include/bits/sys_errlist.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/sys_errlist.h cuobjdumpInstList.h
+cuobjdump_to_ptxplus.o: cuobjdumpInst.h /usr/include/string.h
+cuobjdump_to_ptxplus.o: /usr/include/xlocale.h /usr/include/stdlib.h
+cuobjdump_to_ptxplus.o: /usr/include/sys/types.h /usr/include/time.h
+cuobjdump_to_ptxplus.o: /usr/include/endian.h /usr/include/bits/endian.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/byteswap.h
+cuobjdump_to_ptxplus.o: /usr/include/sys/select.h /usr/include/bits/select.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/sigset.h /usr/include/bits/time.h
+cuobjdump_to_ptxplus.o: /usr/include/sys/sysmacros.h
+cuobjdump_to_ptxplus.o: /usr/include/bits/pthreadtypes.h
+cuobjdump_to_ptxplus.o: /usr/include/alloca.h stringList.h
+elf_lexer.o: /usr/include/stdio.h /usr/include/features.h
+elf_lexer.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+elf_lexer.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+elf_lexer.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+elf_lexer.o: /usr/include/libio.h /usr/include/_G_config.h
+elf_lexer.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+elf_lexer.o: /usr/include/bits/sys_errlist.h /usr/include/string.h
+elf_lexer.o: /usr/include/xlocale.h /usr/include/errno.h
+elf_lexer.o: /usr/include/bits/errno.h /usr/include/linux/errno.h
+elf_lexer.o: /usr/include/asm/errno.h /usr/include/asm-generic/errno.h
+elf_lexer.o: /usr/include/asm-generic/errno-base.h /usr/include/stdlib.h
+elf_lexer.o: /usr/include/sys/types.h /usr/include/time.h
+elf_lexer.o: /usr/include/endian.h /usr/include/bits/endian.h
+elf_lexer.o: /usr/include/bits/byteswap.h /usr/include/sys/select.h
+elf_lexer.o: /usr/include/bits/select.h /usr/include/bits/sigset.h
+elf_lexer.o: /usr/include/bits/time.h /usr/include/sys/sysmacros.h
+elf_lexer.o: /usr/include/bits/pthreadtypes.h /usr/include/alloca.h
+elf_lexer.o: elf_parser.hh /usr/include/unistd.h
+elf_lexer.o: /usr/include/bits/posix_opt.h /usr/include/bits/confname.h
+elf_lexer.o: /usr/include/getopt.h
+elf_parser.o: /usr/include/stdio.h /usr/include/features.h
+elf_parser.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+elf_parser.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+elf_parser.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+elf_parser.o: /usr/include/libio.h /usr/include/_G_config.h
+elf_parser.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+elf_parser.o: /usr/include/bits/sys_errlist.h cuobjdumpInstList.h
+elf_parser.o: cuobjdumpInst.h /usr/include/string.h /usr/include/xlocale.h
+elf_parser.o: /usr/include/stdlib.h /usr/include/sys/types.h
+elf_parser.o: /usr/include/time.h /usr/include/endian.h
+elf_parser.o: /usr/include/bits/endian.h /usr/include/bits/byteswap.h
+elf_parser.o: /usr/include/sys/select.h /usr/include/bits/select.h
+elf_parser.o: /usr/include/bits/sigset.h /usr/include/bits/time.h
+elf_parser.o: /usr/include/sys/sysmacros.h /usr/include/bits/pthreadtypes.h
+elf_parser.o: /usr/include/alloca.h stringList.h
+stringList.o: stringList.h /usr/include/stdio.h /usr/include/features.h
+stringList.o: /usr/include/sys/cdefs.h /usr/include/bits/wordsize.h
+stringList.o: /usr/include/gnu/stubs.h /usr/include/gnu/stubs-64.h
+stringList.o: /usr/include/bits/types.h /usr/include/bits/typesizes.h
+stringList.o: /usr/include/libio.h /usr/include/_G_config.h
+stringList.o: /usr/include/wchar.h /usr/include/bits/stdio_lim.h
+stringList.o: /usr/include/bits/sys_errlist.h
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.cc b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
index 524c0da..e0d5557 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.cc
@@ -1,6 +1,10 @@
-#include "cuobjdumpInst.h"
+// External includes
#include <sstream>
#include <iostream>
+#include <cassert>
+
+// Local includes
+#include "cuobjdumpInst.h"
extern void output(const char * text);
@@ -8,6 +12,191 @@ extern void output(const char * text);
cuobjdumpInst::cuobjdumpInst()
{
//initilize everything to empty
+ m_label = "";
+ m_predicate = new stringList();
+ m_base = "";
+ m_baseModifiers = new stringList();
+ m_typeModifiers = new stringList();
+ m_operands = new stringList();
+ m_predicateModifiers = new stringList();
+}
+
+void cuobjdumpInst::printCuobjdumpInst()
+{
+ /*TODO: print label here*/
+ /*TODO: print predicate here*/
+ /*std::cout << "Instruction Base: " << m_base << "\n";
+
+ std::cout << "Instruction Modifiers: ";
+ m_baseModifiers->printStringList();
+ std::cout << "\n";
+
+ std::cout << "Operand types: ";
+ m_typeModifiers->printStringList();
+ std::cout << "\n";
+
+ std::cout << "Operands: ";
+ m_operands->printStringList();
+ std::cout << "\n\n";*/
+
+ std::cout << m_base << " ";
+ m_baseModifiers->printStringList();
+ std::cout << " ";
+ m_typeModifiers->printStringList();
+ std::cout << " ";
+ m_operands->printStringList();
+ std::cout << "\n";
+}
+
+// Just prints the base and operands
+void cuobjdumpInst::printHeaderPtx()
+{
+ output(m_base); output(" ");
+
+ stringListPiece* currentPiece;
+
+ currentPiece = m_baseModifiers->getListStart();
+ for(int i=0; (i<m_baseModifiers->getSize())&&(currentPiece!=NULL); i++)
+ {
+ output(" "); output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+ }
+
+ currentPiece = m_operands->getListStart();
+ for(int i=0; (i<m_operands->getSize())&&(currentPiece!=NULL); i++)
+ {
+ output(" "); output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+ }
+}
+
+//retreive instruction mnemonic
+const char* cuobjdumpInst::getBase()
+{
+ return m_base;
+}
+
+stringList* cuobjdumpInst::getTypeModifiers()
+{
+ return m_typeModifiers;
+}
+
+//print out .version and .target header lines
+bool cuobjdumpInst::printHeaderInst()
+{
+ if(strcmp(m_base, ".version")==0)
+ {
+ output(m_base); output(" ");
+
+ stringListPiece* currentPiece = m_operands->getListStart();
+ output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+
+ if(currentPiece!=NULL)
+ {
+ output("."); output(currentPiece->stringText);
+ }
+ output("+");
+ output("\n");
+ }
+ else if(strcmp(m_base, ".target")==0)
+ {
+ output(m_base); output(" ");
+
+ stringListPiece* currentPiece = m_operands->getListStart();
+ output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+
+ while(currentPiece!=NULL)
+ {
+ output(", "); output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+ }
+ output("\n");
+ }
+ else if(strcmp(m_base, ".tex")==0)
+ {
+ output(m_base); output(" ");
+
+ stringListPiece* currentPiece;
+
+
+ currentPiece = m_baseModifiers->getListStart();
+ output(currentPiece->stringText); output(" ");
+ currentPiece = currentPiece->nextString;
+
+ while(currentPiece!=NULL)
+ {
+ output(" "); output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+ }
+
+ currentPiece = m_operands->getListStart();
+ output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+
+ while(currentPiece!=NULL)
+ {
+ output(" "); output(currentPiece->stringText);
+ currentPiece = currentPiece->nextString;
+ }
+ output(";\n");
+ }
+ else
+ {
+ return false;
+ }
+ return true;
+}
+
+void cuobjdumpInst::setBase(const char* setBaseValue)
+{
+ m_base = setBaseValue;
+}
+
+void cuobjdumpInst::addBaseModifier(const char* addBaseMod)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addBaseMod;
+
+ m_baseModifiers->add(tempPiece);
+}
+
+void cuobjdumpInst::addTypeModifier(const char* addTypeMod)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addTypeMod;
+
+ m_typeModifiers->add(tempPiece);
+}
+
+void cuobjdumpInst::addOperand(const char* addOp)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addOp;
+
+ m_operands->add(tempPiece);
+}
+
+void cuobjdumpInst::setPredicate(const char* setPredicateValue)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = setPredicateValue;
+
+ m_predicate->add(tempPiece);
+}
+
+void cuobjdumpInst::addPredicateModifier(const char* addPredicateMod)
+{
+ stringListPiece* tempPiece = new stringListPiece;
+ tempPiece->stringText = addPredicateMod;
+
+ m_predicateModifiers->add(tempPiece);
+}
+
+void cuobjdumpInst::setLabel(const char* setLabelValue)
+{
+ m_label = setLabelValue;
}
bool cuobjdumpInst::checkCubojdumpLabel(std::list<std::string> labelList, std::string label)
@@ -182,7 +371,7 @@ void cuobjdumpInst::printCuobjdumpBaseModifiers()
}
else if( strcmp(modString, ".s")==0 )
{
- //".s" is the same as ".join" in decuda.
+ //".s" is the same as ".join" in cuobjdump.
}
else if( strcmp(modString, ".sfu")==0 )
{
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInst.h b/cuobjdump_to_ptxplus/cuobjdumpInst.h
index 20d31be..8cf0ebe 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInst.h
+++ b/cuobjdump_to_ptxplus/cuobjdumpInst.h
@@ -1,12 +1,45 @@
-#include "decudaInstList.h"
+#ifndef _CUOBJDUMPINST_H_
+#define _CUOBJDUMPINST_H_
-class cuobjdumpInst: public decudaInst
+// External includes
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <list>
+
+// Local includes
+//#include "cuobjdumpInstList.h"
+#include "stringList.h"
+
+class cuobjdumpInst
{
+protected:
+ //instruction data
+ const char* m_label; //instruction label
+ stringList* m_predicate; //instruction predicate
+ const char* m_base; //instruction mnemonic
+ stringList* m_baseModifiers; //base modifiers
+ stringList* m_typeModifiers; //operand types
+ stringList* m_operands; //operands
+ stringList* m_predicateModifiers; //predicate modifiers
public:
//Constructor
cuobjdumpInst();
+ //accessors
+ const char* getBase();
+ stringList* getTypeModifiers();
+
+ //Mutators
+ void setLabel(const char* setLabelValue);
+ void setPredicate(const char* setPredicateValue);
+ void addPredicateModifier(const char* addPredicateMod);
+ void setBase(const char* setBaseValue);
+ void addBaseModifier(const char* addBaseMod);
+ void addTypeModifier(const char* addTypeMod);
+ void addOperand(const char* addOp);
+
bool checkCubojdumpLabel(std::list<std::string> labelList, std::string label);
void printCuobjdumpLabel(std::list<std::string> labelList);
@@ -18,5 +51,11 @@ public:
void printCuobjdumpOperands();
void printCuobjdumpPtxPlus(std::list<std::string> labelList, std::list<std::string> texList);
+
+ //print representation
+ bool printHeaderInst();
+ void printCuobjdumpInst();
+ void printHeaderPtx();
};
+#endif //_CUOBJDUMPINST_H_
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
index 5eaf67f..45c8f9d 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.cc
@@ -1,6 +1,17 @@
+#include <sstream>
#include <iostream>
+#include <cassert>
#include "cuobjdumpInstList.h"
+#define P_DEBUG 1
+#define DPRINTF(...) \
+ if(P_DEBUG) { \
+ printf("(%s:%u) ", __FILE__, __LINE__); \
+ printf(__VA_ARGS__); \
+ printf("\n"); \
+ fflush(stdout); \
+ }
+
extern void output(const char * text);
//Constructor
@@ -11,6 +22,289 @@ cuobjdumpInstList::cuobjdumpInstList()
+// add to tex list
+void cuobjdumpInstList::addTex(std::string tex)
+{
+ std::string origTex = tex;
+ DPRINTF("cuobjdumpInstList::addTex tex=%s", tex.c_str());
+ // If $tex# tex from cuobjdump, then use index to get real tex name
+ if(tex.substr(0, 4) == "$tex") {
+ tex = tex.substr(4, tex.size()-4);
+ int texNum = atoi(tex.c_str());
+ if(texNum >= m_realTexList.size()) {
+ output("ERROR: tex does not exist in real tex list from ptx.\n.");
+ assert(0);
+ }
+
+ std::list<std::string>::iterator itex = m_realTexList.begin();
+ for(int i=0; i<texNum; i++) itex++;
+ origTex = *itex;
+ }
+ // Otherwise, tex from original ptx
+ else {
+ m_realTexList.push_back(tex);
+ }
+
+ // Add the tex to instruction operand list
+ //char* texName = new char [strlen(origTex.c_str())+1];
+ //strcpy(texName, origTex.c_str());
+ //getListEnd().addOperand(texName);
+}
+
+void cuobjdumpInstList::setLastEntryName(std::string entryName)
+{
+ m_entryList.back().m_entryName = entryName;
+}
+
+// create new global constant memory "bank"
+void cuobjdumpInstList::addConstMemory(int index)
+{
+ constMemory newConstMem;
+ newConstMem.index = index;
+ newConstMem.entryIndex = 0;
+ m_constMemoryList.push_back(newConstMem);
+}
+
+//add cuobjdumpInst to the last entry in entry list
+int cuobjdumpInstList::add(cuobjdumpInst* newCuobjdumpInst)
+{
+ if(m_entryList.size() == 0) {
+ //output("ERROR: Adding an instruction before entry.\n");
+ addEntry("");
+ //assert(0);
+ }
+
+ m_entryList.back().m_instList.push_back(*newCuobjdumpInst);
+
+ return m_entryList.size();
+}
+
+// add a new entry
+int cuobjdumpInstList::addEntry(std::string entryName)
+{
+ cuobjdumpEntry newEntry;
+ newEntry.m_largestRegIndex = -1;
+ newEntry.m_largestOfsRegIndex = -1;
+ newEntry.m_largestPredIndex = -1;
+ newEntry.m_reg124 = false;
+ newEntry.m_oreg127 = false;
+ newEntry.m_lMemSize = -1;
+
+ newEntry.m_entryName = entryName;
+
+
+ // Fill opPerCycle histogram with values
+ newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_1", 0) );
+ newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_2", 0) );
+ newEntry.m_opPerCycleHistogram.insert( std::pair<std::string,int>("OP_8", 0) );
+
+
+ m_entryList.push_back(newEntry);
+ return m_entryList.size();
+}
+
+// print out .version and .target headers
+void cuobjdumpInstList::printHeaderInstList()
+{
+ // These should be in the first entry
+ cuobjdumpEntry e_first = m_entryList.front();
+
+ std::list<cuobjdumpInst>::iterator currentInst;
+ for(currentInst=e_first.m_instList.begin(); currentInst!=e_first.m_instList.end(); ++currentInst)
+ {
+ if(!(currentInst->printHeaderInst()))
+ {
+ break;
+ }
+ }
+ for ( std::list<std::string>::iterator iter = m_realTexList.begin();
+ iter != m_realTexList.end();
+ iter ++) {
+ output(".tex .u64 ");
+ output((*iter).c_str());
+ output(";\n");
+ }
+}
+
+bool cuobjdumpInstList::findEntry(std::string entryName, cuobjdumpEntry& entry) {
+ std::list<cuobjdumpEntry>::iterator e;
+
+ std::string entryNameS = entryName;
+
+ for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
+ if( e->m_entryName == entryNameS) {
+ entry = *e;
+ return true;
+ }
+ }
+
+ return false;
+}
+
+// get the list of real tex names
+std::list<std::string> cuobjdumpInstList::getRealTexList() {
+ return m_realTexList;
+}
+
+// set the list of real tex names
+void cuobjdumpInstList::setRealTexList(std::list<std::string> realTexList) {
+ m_realTexList = realTexList;
+}
+
+// add value to const memory
+void cuobjdumpInstList::addConstMemoryValue(std::string constMemoryValue)
+{
+ m_constMemoryList.back().m_constMemory.push_back(constMemoryValue);
+}
+
+void cuobjdumpInstList::addConstMemoryValue2(std::string constMemoryValue)
+{
+ m_constMemoryList2.back().m_constMemory.push_back(constMemoryValue);
+}
+
+// set type of constant memory
+void cuobjdumpInstList::setConstMemoryType(const char* type)
+{
+ m_constMemoryList.back().type = type;
+}
+
+void cuobjdumpInstList::setConstMemoryType2(const char* type)
+{
+ m_constMemoryList2.back().type = type;
+}
+
+//retrieve point to list end
+cuobjdumpInst cuobjdumpInstList::getListEnd()
+{
+ return m_entryList.back().m_instList.back();
+}
+
+// print out predicate names
+void cuobjdumpInstList::printPredNames(cuobjdumpEntry entry)
+{
+ if( entry.m_largestPredIndex >= 0) {
+ char out[30];
+ sprintf(out, "\t.reg .pred $p<%d>;", entry.m_largestPredIndex+1);
+ output(out);
+ output("\n");
+ }
+
+}
+
+// print reg124 and set its value to 0
+void cuobjdumpInstList::printOutOfBoundRegisters(cuobjdumpEntry entry)
+{
+ if( entry.m_reg124 == true ) {
+ output("\n");
+ output("\t.reg .u32 $r124;\n");
+ output("\tmov.u32 $r124, 0x00000000;\n");
+ }
+ if( entry.m_oreg127 == true) {
+ output("\n");
+ output("\t.reg .u32 $o127;\n");
+ }
+}
+
+// print out register names
+void cuobjdumpInstList::printRegNames(cuobjdumpEntry entry)
+{
+ if( entry.m_largestRegIndex >= 0) {
+ char out[30];
+ sprintf(out, "\t.reg .u32 $r<%d>;", entry.m_largestRegIndex+1);
+ output(out);
+ output("\n");
+ }
+
+ if( entry.m_largestOfsRegIndex >= 0) {
+ char out[30];
+ sprintf(out, "\t.reg .u32 $ofs<%d>;", entry.m_largestOfsRegIndex+1);
+ output(out);
+ output("\n");
+ }
+}
+
+// print const memory directive
+void cuobjdumpInstList::printMemory()
+{
+
+ // Constant memory
+
+ for(std::list<constMemory>::iterator i=m_constMemoryList.begin(); i!=m_constMemoryList.end(); ++i) {
+ char line[40];
+
+ // Global or entry specific
+ if(i->entryIndex == 0)
+ sprintf(line, ".const %s constant0[%d] = {", i->type, i->m_constMemory.size());
+ else
+ sprintf(line, ".const %s ce%dc%d[%d] = {", i->type, i->entryIndex, i->index, i->m_constMemory.size());
+
+ output(line);
+
+ std::list<std::string>::iterator j;
+ int l=0;
+ for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
+ if(j!=i->m_constMemory.begin())
+ output(", ");
+ if( (l++ % 4) == 0) output("\n ");
+ output(j->c_str());
+ }
+ output("\n};\n\n");
+ }
+
+
+ for(std::list<constMemory2>::iterator i=m_constMemoryList2.begin(); i!=m_constMemoryList2.end(); ++i) {
+ char line[1024];
+
+ // Global or entry specific
+ sprintf(line, ".const %s constant1%s[%d] = {", i->type, i->kernel, i->m_constMemory.size());
+
+ output(line);
+
+ std::list<std::string>::iterator j;
+ int l=0;
+ for(j=i->m_constMemory.begin(); j!=i->m_constMemory.end(); ++j) {
+ if(j!=i->m_constMemory.begin())
+ output(", ");
+ if( (l++ % 4) == 0) output("\n ");
+ output(j->c_str());
+ }
+ output("\n};\n\n");
+ }
+
+ // Next, print out the local memory declaration
+ std::list<cuobjdumpEntry>::iterator e;
+ int eIndex=1; // entry index starts from 1 from the first blank entry is missing here (only in header entry list)
+ for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
+ if(e->m_lMemSize > 0) {
+ std::stringstream ssout;
+ ssout << ".local .b8 l" << eIndex << "[" << e->m_lMemSize << "];" << std::endl;
+ output(ssout.str().c_str());
+ }
+ eIndex++;
+ }
+ output("\n");
+
+ // Next, print out the global memory declaration
+ std::list<globalMemory>::iterator g;
+ for(g=m_globalMemoryList.begin(); g!=m_globalMemoryList.end(); ++g) {
+ std::stringstream out;
+ out << ".global .b8 " << g->name << "[" << g->bytes << "];" << std::endl;
+ output(out.str().c_str());
+ }
+ output("\n");
+
+ // Next, print out constant memory pointers
+ std::list<constMemoryPtr>::iterator cp;
+ for(cp=m_constMemoryPtrList.begin(); cp!=m_constMemoryPtrList.end(); ++cp) {
+ std::stringstream out;
+ out << ".const .b8 " << cp->name << "[" << cp->bytes << "];" << std::endl;
+ out << ".constptr " << cp->name << ", " << cp->destination << ", " << cp->offset << ";" << std::endl;
+ output(out.str().c_str());
+ }
+ output("\n");
+
+}
+
//TODO: Some register processing work is supposed to be done here.
void cuobjdumpInstList::addCuobjdumpRegister(std::string reg, bool lo)
@@ -303,164 +597,25 @@ void cuobjdumpInstList::addEntryLocalMemory(int value, int entryIndex)
m_localMemoryList.push_back(newLocalMem);
}
-
-// Read in constant memory from bin file
-// Two cases of constant memory have been noticed so far
-// 1 - All the constant memory is initialized in original ptx file. The assembler combines all this memory into c0
-// 2 - Constant memory is declared in ptx, but not initialized (initialized by host). The assembler still calls this c0
-void cuobjdumpInstList::readConstMemoryFromElfFile(std::string elf)
-{
- unsigned k=1;
- printf("Trying to find constant memory in elf file:\n");
-
- // Get each constant segment
- const boost::regex constPattern("^\\.nv\\.constant1\\.[^\n]+\n[ x0-9a-f\t]+$");
- // Parse each constseg
- const boost::sregex_token_iterator end;
- for (
- boost::sregex_token_iterator i(elf.begin(),elf.end(), constPattern);
- i != end;
- ++i
- )
- {
- std::string memseg = *i;
- boost::smatch memResult;
-
- const boost::regex memValuePattern("(0x[A-Fa-f0-9]{8,8})");
-
- bool memExists = boost::regex_search(memseg, memResult, memValuePattern);
-
-
- std::list<std::string> c1;
- std::list<std::string>::iterator it = c1.begin();
-
- const boost::sregex_token_iterator end2;
- for (
- boost::sregex_token_iterator j(memseg.begin(),memseg.end(), memValuePattern);
- j != end2;
- ++j ){
- c1.insert(it, *j);
- }
-
- addEntryConstMemory(1, k);
- setConstMemoryType(".u32");
-
- std::list<std::string>::iterator c;
- if(c1.size() > 0) {
- for(c=c1.begin(); c!=c1.end(); ++c) {
- std::string a = *c;
- //printf("%s ", a.c_str());
- addConstMemoryValue(a);
- }
- }
-
-
- printf("Found constant memory\n");
- printf(memseg.c_str());
- printf("\n");
- k++;
- }
- m_kernelCount = k-1;
-}
void cuobjdumpInstList::setKernelCount(int k){
m_kernelCount = k;
}
-void cuobjdumpInstList::readOtherConstMemoryFromBinFile(std::string binString)
-{
- // Initialize a list to store memory values
- // std::list<std::string> c0;
-
- // Get each code segment
- //const boost::regex codePattern("(code \\{[^\\{\\}]*(const \\{[^\\{\\}]*(mem \\{[^\\{\\}]*\\}[^\\{\\}]*)+\\}[^\\{\\}]*)+bincode \\{[^\\{\\}]*\\}[^\\{\\}]*\\})");
- const boost::regex codePattern("(code \\{[^\\{\\}]*(const \\{[^\\{\\}]*(mem \\{[^\\{\\}]*\\}[^\\{\\}]*)+\\}[^\\{\\}]*)*bincode \\{[^\\{\\}]*\\}[^\\{\\}]*\\})");
-
- int k=1;
-
- // Parse each codeseg
- const boost::sregex_token_iterator end;
- for(
- boost::sregex_token_iterator i(binString.begin(),binString.end(), codePattern);
- i != end;
- ++i
- )
- {
- std::list<std::string> c1;
-
- // For each code segment, get the seg numbers and memory values string
- std::string codeSeg_s = *i;
- std::string segnum_s, lmem_s, mem;
- int segnum;
- int lmem;
-
- boost::smatch segnumResult;
- boost::smatch lmemResult;
- boost::smatch memResult;
-
- const boost::regex segnumPattern("segnum\\s*=\\s(\\d*)");
- const boost::regex lmemPattern("lmem\\s*=\\s(\\d*)");
- const boost::regex memPattern("mem \\{([^\\}]*)\\}");
-
- boost::regex_search(codeSeg_s, segnumResult, segnumPattern);
- boost::regex_search(codeSeg_s, lmemResult, lmemPattern);
- bool memExists = boost::regex_search(codeSeg_s, memResult, memPattern);
-
- lmem_s = lmemResult[1];
- lmem = atoi(lmem_s.c_str());
-
- addEntryLocalMemory(lmem, k);
-
- if(memExists)
- {
- segnum_s = segnumResult[1];
- segnum = atoi(segnum_s.c_str());
-
- mem = memResult[1];
- const boost::regex memValuePattern("(0x[A-Fa-f0-9]{8,8})");
-
- std::list<std::string>::iterator it = c1.begin();
-
- const boost::sregex_token_iterator end2;
- for (
- boost::sregex_token_iterator j(mem.begin(),mem.end(), memValuePattern);
- j != end2;
- ++j
- )
- {
- c1.insert(it, *j);
- }
-
- addEntryConstMemory(segnum, k);
- setConstMemoryType(".u32");
-
- std::list<std::string>::iterator c;
- if(c1.size() > 0) {
- for(c=c1.begin(); c!=c1.end(); ++c) {
- std::string a = *c;
- //printf("%s ", a.c_str());
- addConstMemoryValue(a);
- }
- }
- }
- k++;
- }
- m_kernelCount = k-1;
-}
void cuobjdumpInstList::printCuobjdumpInstList()
{
// Each entry
- std::list<decudaEntry>::iterator e;
+ std::list<cuobjdumpEntry>::iterator e;
for(e=m_entryList.begin(); e!=m_entryList.end(); ++e) {
- for( std::list<decudaInst>::iterator currentInst=e->m_instList.begin();
+ for( std::list<cuobjdumpInst>::iterator currentInst=e->m_instList.begin();
currentInst!=e->m_instList.end();
++currentInst) {
// Output the instruction
output("\t");
- currentInst->printDecudaInst();
+ currentInst->printCuobjdumpInst();
output("\n");
}
}
@@ -486,18 +641,18 @@ void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
printMemory();
printCuobjdumpLocalMemory();
// Each entry
- std::list<decudaEntry>::reverse_iterator e;
+ std::list<cuobjdumpEntry>::reverse_iterator e;
for(e=m_entryList.rbegin(); e!=m_entryList.rend(); ++e) {
output("\n");
// Output the header information for this entry using headerInfo
// First, find the matching entry in headerInfo
- decudaEntry headerEntry;
+ cuobjdumpEntry headerEntry;
if( headerInfo->findEntry(e->m_entryName, headerEntry) ) {
// Entry for current header found, print it out
- std::list<decudaInst>::iterator headerInstIter;
+ std::list<cuobjdumpInst>::iterator headerInstIter;
for(headerInstIter=headerEntry.m_instList.begin();
headerInstIter!=headerEntry.m_instList.end();
++headerInstIter) {
@@ -517,7 +672,7 @@ void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
output("\n");
output("{\n");
} else {
- output("Mismatch in entry names between decuda output and original ptx file.\n");
+ output("Mismatch in entry names between cuobjdump output and original ptx file.\n");
assert(0);
}
}
@@ -527,13 +682,12 @@ void cuobjdumpInstList::printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo)
printOutOfBoundRegisters(*e);
output("\n");
- for(std::list<decudaInst>::iterator currentInst=e->m_instList.begin(); currentInst!=e->m_instList.end(); ++currentInst){
+ for(std::list<cuobjdumpInst>::iterator currentInst=e->m_instList.begin(); currentInst!=e->m_instList.end(); ++currentInst){
// Output the instruction
//cuobjdumpInst* outputInst = &*currentInst;
- cuobjdumpInst* outputInst = static_cast<cuobjdumpInst*>(&*currentInst);
output("\t");
//outputInst->printCuobjdumpPtxPlus(m_entryList.back().m_labelList);
- outputInst->printCuobjdumpPtxPlus(e->m_labelList, this->m_realTexList);
+ currentInst->printCuobjdumpPtxPlus(e->m_labelList, this->m_realTexList);
output("\n");
}
output("\n\tl_exit: exit;\n");
diff --git a/cuobjdump_to_ptxplus/cuobjdumpInstList.h b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
index 1bc8f10..7cb3a32 100644
--- a/cuobjdump_to_ptxplus/cuobjdumpInstList.h
+++ b/cuobjdump_to_ptxplus/cuobjdumpInstList.h
@@ -1,5 +1,73 @@
+#ifndef _CUOBJDUMPINSTLIST_H_
+#define _CUOBJDUMPINSTLIST_H_
+
+// External includes
+#include <list>
+#include <map>
+#include <string>
+
+// Local includes
#include "cuobjdumpInst.h"
+// Used for entry specific constant memory segments (c1)
+struct constMemory
+{
+ int index;
+ int entryIndex;
+ const char* type;
+ std::list<std::string> m_constMemory;
+};
+
+struct constMemory2
+{
+ const char* kernel;
+ const char* type;
+ std::list<std::string> m_constMemory;
+};
+
+// Used for uninitialized constant memory (globally defined)
+struct constMemoryPtr
+{
+ int bytes;
+ std::string name;
+
+ std::string destination;
+ int offset;
+};
+
+// Used for global memory segments
+struct globalMemory
+{
+ int offset;
+ int bytes;
+ std::string name;
+};
+
+struct cuobjdumpEntry
+{
+ //char* m_entryName;
+ std::string m_entryName;
+ std::list<cuobjdumpInst> m_instList; // List of cuobjdump instructions
+
+ // Register list
+ int m_largestRegIndex;
+ int m_largestOfsRegIndex;
+ bool m_reg124;
+ bool m_oreg127;
+
+ // Predicate list
+ int m_largestPredIndex;
+
+ // Local memory size
+ int m_lMemSize;
+
+ //use for recording used labels
+ std::list<std::string> m_labelList;
+
+ // Histogram for operation per cycle count
+ std::map<std::string, int> m_opPerCycleHistogram;
+};
+
// Used for local memory segments
struct localMemory
{
@@ -7,22 +75,46 @@ struct localMemory
int entryIndex;
};
-class cuobjdumpInstList: public decudaInstList
+class cuobjdumpInstList
{
protected:
- std::string parseCuobjdumpPredicate(std::string pred);
+ std::list<cuobjdumpEntry> m_entryList;
+ std::list<constMemory> m_constMemoryList;
+ std::list<constMemory2> m_constMemoryList2;
+ std::list<globalMemory> m_globalMemoryList;
+
int m_kernelCount;
std::map<std::string,int>kernelcmemmap;
std::map<std::string,int>kernellmemmap;
std::list<localMemory> m_localMemoryList;
+ std::list<std::string> m_realTexList; // Stores the real names of tex variables
+ std::list<constMemoryPtr> m_constMemoryPtrList;
+
+ // Functions:
+ std::string parseCuobjdumpPredicate(std::string pred);
+ void printMemory();// Print const memory directives
+ // Print register names
+ void printRegNames(cuobjdumpEntry entry);
+ void printOutOfBoundRegisters(cuobjdumpEntry entry);
+ // Print predicate names
+ void printPredNames(cuobjdumpEntry entry);
public:
//Constructor
cuobjdumpInstList();
+ cuobjdumpInst getListEnd();
+
+ // Functions used by the parser
+ int addEntry(std::string entryName); // creates a new entry
+ int add(cuobjdumpInst* newInst); //add cuobjdumpInst to list
+ void addConstMemory(int index); // add global const memory
+ void addTex(std::string tex); // add tex operand
+ bool findEntry(std::string entryName, cuobjdumpEntry& entry); // find and return entry
+
void setKernelCount(int k);
void readConstMemoryFromElfFile(std::string elf);
-
+ void setLastEntryName(std::string entryName); // sets name of last entry
void addCuobjdumpRegister(std::string reg, bool lo=false); //add register
void addCuobjdumpMemoryOperand(std::string mem, int memType);
std::string parseCuobjdumpRegister(std::string reg, bool lo, int vectorFlag);
@@ -32,16 +124,22 @@ public:
void addEntryConstMemory(int index, int entryIndex);
void addEntryConstMemory2(char* kernel);
+ void setConstMemoryType(const char* type);
+ void setConstMemoryType2(const char* type); // set type of constant memory
+ void addConstMemoryValue(std::string constMemoryValue); // add const memory
+ void addConstMemoryValue2(std::string constMemoryValue); // add const memory
void addConstMemoryPtr(const char* bytes, const char* offset, const char* name);
void setConstMemoryMap(const char* kernelname, int index);
void setLocalMemoryMap(const char* kernelname, int index);
void reverseConstMemory();
void addEntryLocalMemory(int value, int entryIndex);
void readOtherConstMemoryFromBinFile(std::string binString); // read in constant memory from bin file
-
+ std::list<std::string> getRealTexList(); // get the list of real tex names
+ void setRealTexList(std::list<std::string> realTexList); // set the list of real tex names
+ void printHeaderInstList();
void printCuobjdumpLocalMemory();
-
void printCuobjdumpInstList();
void printCuobjdumpPtxPlusList(cuobjdumpInstList* headerInfo);
};
+#endif //_CUOBJDUMPINSTLIST_H_
diff --git a/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc b/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
index e8d45ef..fa322e1 100644
--- a/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
+++ b/cuobjdump_to_ptxplus/cuobjdump_to_ptxplus.cc
@@ -1,7 +1,9 @@
#include <iostream>
-#include "cuobjdumpInstList.h"
#include <stdio.h>
#include <fstream>
+#include <cassert>
+
+#include "cuobjdumpInstList.h"
using namespace std;
diff --git a/cuobjdump_to_ptxplus/elf.l b/cuobjdump_to_ptxplus/elf.l
index 3457a7a..f594a00 100644
--- a/cuobjdump_to_ptxplus/elf.l
+++ b/cuobjdump_to_ptxplus/elf.l
@@ -2,7 +2,7 @@
%{
#include <stdio.h>
#include <string.h>
-#include "elf_parser.hpp"
+#include "elf_parser.hh"
#define yylval elf_lval
diff --git a/cuobjdump_to_ptxplus/header.l b/cuobjdump_to_ptxplus/header.l
new file mode 100644
index 0000000..170ecf7
--- /dev/null
+++ b/cuobjdump_to_ptxplus/header.l
@@ -0,0 +1,124 @@
+/*
+// Copyright (c) 2009-2011, Jimmy Kwa,
+// The University of British Columbia
+// 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 The University of British Columbia 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 AND CONTRIBUTORS "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 HOLDER 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.
+*/
+
+%option noyywrap
+%option yylineno
+%option prefix="header_"
+%{
+#include <iostream>
+#include "header_parser.hh"
+#include <string.h>
+
+/* not sure if this is needed*/
+#define YY_NEVER_INTERACTIVE 1
+
+void header_error(const char*);
+extern void output(const char * text);
+%}
+
+
+/*Regular expresions go here*/
+alpha [a-zA-Z]
+numeric [0-9]
+hexalpha [A-Fa-f]
+alphanum ({alpha}|{numeric})
+hexdigit ({numeric}|{hexalpha})
+whitespace [ \t]
+newline [\n]
+anycharbutnewline [^\n]
+anychar ({newline}|{anycharbutnewline})
+anycharbutrightbrace [^"}"]
+
+identifier ({alpha}({alphanum}|["_$"])*)|([_$%]({alphanum}|["_$"])+)
+hexliteral 0[xX]{hexdigit}+U?
+octliteral 0[0-7]+U?
+binliteral 0[bB][01]+U?
+decliteral ([1-9][0-9]*U?)|0U?
+
+%%
+ /*Translation rules*/
+
+ /*Compiler directives*/
+\.version return DOTVERSION;
+\.target return DOTTARGET;
+\.entry return DOTENTRY;
+
+ /*State spaces*/
+\.param return DOTPARAM;
+
+ /*operand types*/
+\.u64 return DOTU64;
+\.u32 return DOTU32;
+\.u16 return DOTU16;
+\.b32 return DOTB32;
+\.f32 return DOTF32;
+
+
+ /*identifier + literals*/
+{identifier} header_lval.string_value = strdup(yytext); return IDENTIFER;
+{decliteral} header_lval.string_value = strdup(yytext); return DECLITERAL; //change to int later?
+
+ /*braces and parenthesis*/
+\( return LEFTPAREN;
+\) return RIGHTPAREN;
+
+ /*discard newline*/
+{newline}+
+
+
+ /*discard commas and periods*/
+\,
+\.
+ /*discard whitespace*/
+{whitespace}+
+ /*discard comments*/
+"//"{anycharbutnewline}*{newline}
+
+ /*discard other header stuff*/
+\.file{anycharbutnewline}*{newline}
+\.extern{anycharbutnewline}*{newline}
+
+ /*discard kernel body*/
+\{{anycharbutrightbrace}*\}
+
+ /*reached end of file*/
+<<EOF>> return 0;
+
+ /*None of the rules matched. Throw error*/
+. header_error("Invalid token");
+
+%%
+
+void header_error(const char* message)
+{
+ output(" "); output(message); output(" near \""); output(header_text); output("\"");
+ output(" on line ");
+ char line[5]; sprintf(line, "%i", header_lineno); output(line);
+ output("\n");
+}
diff --git a/cuobjdump_to_ptxplus/header.y b/cuobjdump_to_ptxplus/header.y
new file mode 100644
index 0000000..1bf62ef
--- /dev/null
+++ b/cuobjdump_to_ptxplus/header.y
@@ -0,0 +1,116 @@
+// Copyright (c) 2009-2011, Jimmy Kwa,
+// The University of British Columbia
+// 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 The University of British Columbia 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 AND CONTRIBUTORS "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 HOLDER 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.
+
+
+%{
+#include <iostream>
+#include "cuobjdumpInstList.h"
+
+int yylex(void);
+void yyerror(const char*);
+
+extern cuobjdumpInstList *g_headerList;
+extern void output(const char * text);
+%}
+
+
+%union {
+ double double_value;
+ float float_value;
+ int int_value;
+ char * string_value;
+ void * ptr_value;
+}
+
+%token DOTVERSION
+%token DOTTARGET
+%token DOTENTRY
+
+%token DOTPARAM
+
+%token DOTU64
+%token DOTU32
+%token DOTU16
+%token DOTB32
+%token DOTF32
+
+%token <string_value> IDENTIFER
+
+ /*change these 4 to int later?*/
+%token <string_value> DECLITERAL
+
+%token LEFTPAREN
+%token RIGHTPAREN
+
+
+%%
+
+ /*translation rules*/
+program : statementList { output("No parsing errors\n"); }
+ ;
+
+statementList : statementList statement { output("\n"); }
+ | statement { output("\n"); }
+ ;
+
+statement : compilerDirective literal literal {}
+ | compilerDirective identifierList {}
+ | compilerDirective identifierList LEFTPAREN parameterList RIGHTPAREN {}
+ ;
+
+compilerDirective : DOTVERSION { output(".version"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".version"); g_headerList->add(instEntry);}
+ | DOTTARGET { output(".target"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".target"); g_headerList->add(instEntry);}
+ | DOTENTRY { output(".entry"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".entry"); g_headerList->add(instEntry);}
+ ;
+
+identifierList : identifierList IDENTIFER { output(" "); output($2); g_headerList->getListEnd().addOperand($2); }
+ | IDENTIFER { output(" "); output($1); g_headerList->getListEnd().addOperand($1); }
+ ;
+
+parameterList : parameterList parameter
+ | parameter
+ ;
+
+parameter : stateSpace opTypes IDENTIFER { output(" "); output($3); g_headerList->getListEnd().addOperand($3);}
+ ;
+
+stateSpace : DOTPARAM { output("\n.param"); cuobjdumpInst *instEntry = new cuobjdumpInst(); instEntry->setBase(".param"); g_headerList->add(instEntry); }
+ ;
+
+opTypes : DOTU64 { output(".u64"); g_headerList->getListEnd().addTypeModifier(".u64");}
+ | DOTU32 { output(".u32"); g_headerList->getListEnd().addTypeModifier(".u32");}
+ | DOTU16 { output(".u16"); g_headerList->getListEnd().addTypeModifier(".u16");}
+ | DOTB32 { output(".b32"); g_headerList->getListEnd().addTypeModifier(".b32");}
+ | DOTF32 { output(".f32"); g_headerList->getListEnd().addTypeModifier(".f32");}
+ ;
+
+literal : DECLITERAL { output(" "); output($1); g_headerList->getListEnd().addOperand($1); }
+ ;
+
+%%
+
+/*support c++ functions go here*/
diff --git a/cuobjdump_to_ptxplus/ptx_parser.h b/cuobjdump_to_ptxplus/ptx_parser.h
index 5111743..e32f6d8 100644
--- a/cuobjdump_to_ptxplus/ptx_parser.h
+++ b/cuobjdump_to_ptxplus/ptx_parser.h
@@ -7,8 +7,9 @@
#include <assert.h>
#include <iostream>
#include <sstream>
-#include "decudaInstList.h"
#include <string.h>
+#include <stdio.h>
+#include "cuobjdumpInstList.h"
#define NON_ARRAY_IDENTIFIER 1
#define ARRAY_IDENTIFIER_NO_DIM 2
@@ -79,9 +80,7 @@ void add_constptr(const char* identifier1, const char* identifier2, int offset)
/*non-dummy stuff below this point*/
-
-
-extern decudaInstList *g_headerList;
+extern cuobjdumpInstList *g_headerList;
// Global variable to track if we are currently inside a entry directive
bool inEntryDirective = false;
@@ -118,13 +117,13 @@ void add_function_name( const char *headerInput )
void add_space_spec( enum _memory_space_t spec, int value )
{
DPRINTF("spec=%u", spec);
- decudaInst *instEntry;
+ cuobjdumpInst *instEntry;
static int constmemindex=1;
switch(spec)
{
case param_space_unclassified:
if(inEntryDirective && inParamDirective) {
- instEntry = new decudaInst();
+ instEntry = new cuobjdumpInst();
instEntry->setBase(".param");
g_headerList->add(instEntry);
}
@@ -132,7 +131,7 @@ void add_space_spec( enum _memory_space_t spec, int value )
case tex_space:
inTexDirective = true;
/*
- instEntry = new decudaInst();
+ instEntry = new cuobjdumpInst();
instEntry->setBase(".tex");
g_headerList->add(instEntry);
*/
@@ -141,7 +140,7 @@ void add_space_spec( enum _memory_space_t spec, int value )
if(!inEntryDirective) {
/*
inConstDirective = true;
- instEntry = new decudaInst();
+ instEntry = new cuobjdumpInst();
instEntry->setBase(".const");
g_headerList->add(instEntry);
*/
@@ -219,7 +218,7 @@ void add_scalar_type_spec( int headerInput )
void add_version_info( float versionNumber, unsigned ext)
{
DPRINTF("");
- decudaInst *instEntry = new decudaInst();
+ cuobjdumpInst *instEntry = new cuobjdumpInst();
instEntry->setBase(".version");
g_headerList->add(instEntry);
@@ -236,7 +235,7 @@ void add_version_info( float versionNumber, unsigned ext)
void target_header(char* firstTarget)
{
DPRINTF("%s", firstTarget);
- decudaInst *instEntry = new decudaInst();
+ cuobjdumpInst *instEntry = new cuobjdumpInst();
instEntry->setBase(".target");
g_headerList->add(instEntry);
@@ -246,7 +245,7 @@ void target_header(char* firstTarget)
void target_header2(char* firstTarget, char* secondTarget)
{
DPRINTF("%s, %s", firstTarget, secondTarget);
- decudaInst *instEntry = new decudaInst();
+ cuobjdumpInst *instEntry = new cuobjdumpInst();
instEntry->setBase(".target");
g_headerList->add(instEntry);
@@ -258,7 +257,7 @@ void target_header2(char* firstTarget, char* secondTarget)
void target_header3(char* firstTarget, char* secondTarget, char* thirdTarget)
{
DPRINTF("%s, %s, %s", firstTarget, secondTarget, thirdTarget);
- decudaInst *instEntry = new decudaInst();
+ cuobjdumpInst *instEntry = new cuobjdumpInst();
instEntry->setBase(".target");
g_headerList->add(instEntry);
@@ -290,7 +289,7 @@ void func_header(const char* headerBase)
if((strcmp(headerBase, ".entry")==0)||(strcmp(headerBase, ".func")==0)) {
inEntryDirective = true;
g_headerList->addEntry("");
- decudaInst *instEntry = new decudaInst();
+ cuobjdumpInst *instEntry = new cuobjdumpInst();
instEntry->setBase(headerBase);
g_headerList->add(instEntry);
diff --git a/cuobjdump_to_ptxplus/sass.l b/cuobjdump_to_ptxplus/sass.l
index e665770..859b672 100644
--- a/cuobjdump_to_ptxplus/sass.l
+++ b/cuobjdump_to_ptxplus/sass.l
@@ -1,6 +1,6 @@
%{
#include <stdio.h>
-#include "sass_parser.hpp"
+#include "sass_parser.hh"
#include <string.h>
diff --git a/cuobjdump_to_ptxplus/stringList.cc b/cuobjdump_to_ptxplus/stringList.cc
new file mode 100644
index 0000000..8d1ae42
--- /dev/null
+++ b/cuobjdump_to_ptxplus/stringList.cc
@@ -0,0 +1,128 @@
+// Copyright (c) 2009-2011, Jimmy Kwa,
+// The University of British Columbia
+// 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 The University of British Columbia 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 AND CONTRIBUTORS "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 HOLDER 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.
+
+
+#include "stringList.h"
+#include <stdio.h>
+#include <iostream>
+
+//Constructor
+stringList::stringList()
+{
+ //initilize everything to empty
+ m_size = 0;
+ m_listStart = NULL;
+ m_listEnd = NULL;
+}
+
+//retrieve list size
+int stringList::getSize()
+{
+ return m_size;
+}
+
+//retrieve pointer to list start
+stringListPiece* stringList::getListStart()
+{
+ return m_listStart;
+}
+
+//retrieve point to list end
+stringListPiece* stringList::getListEnd()
+{
+ return m_listEnd;
+}
+
+//add string to end of the list
+//initilize list starting point if this is the first entry
+//increment list size count
+int stringList::add(stringListPiece* newString)
+{
+
+ if(m_listStart==NULL)
+ m_listStart=newString;
+ else
+ m_listEnd->nextString = newString;
+
+ m_listEnd = newString;
+ m_listEnd->nextString = NULL;
+
+ return m_size++;
+}
+
+bool stringList::remove(int index)
+{
+ if(index >= m_size ) return false;
+
+ stringListPiece* m_remove;
+ stringListPiece* currentPiece;
+
+ if(m_size == 1) {
+ m_remove = m_listStart;
+ m_listStart = NULL;
+ m_listEnd = NULL;
+ } else {
+
+ if(index == 0) {
+ m_remove = m_listStart;
+ m_listStart = m_remove->nextString;
+ } else if(index == m_size - 1) {
+ m_remove = m_listEnd;
+ currentPiece = m_listStart;
+ for(int i=1; i<m_size-1; i++)
+ {
+ currentPiece = currentPiece->nextString;
+ }
+ currentPiece->nextString = NULL;
+ m_listEnd = currentPiece;
+ } else {
+ currentPiece = m_listStart;
+ for(int i=1; i<=index-1; i++)
+ {
+ currentPiece = currentPiece->nextString;
+ }
+ m_remove = currentPiece->nextString;
+ currentPiece->nextString = m_remove->nextString;
+ }
+ }
+
+ delete m_remove;
+ m_size -= 1;
+ return true;
+}
+
+//print out all the cuobjdump Instructions in the list
+void stringList::printStringList()
+{
+ stringListPiece* currentPiece = m_listStart;
+
+ for(int i=0; (i<m_size)&&(currentPiece!=NULL); i++)
+ {
+ std::cout << currentPiece->stringText << " ";
+ currentPiece = currentPiece->nextString;
+ }
+}
diff --git a/cuobjdump_to_ptxplus/stringList.h b/cuobjdump_to_ptxplus/stringList.h
new file mode 100644
index 0000000..fe0139e
--- /dev/null
+++ b/cuobjdump_to_ptxplus/stringList.h
@@ -0,0 +1,64 @@
+// Copyright (c) 2009-2011, Jimmy Kwa,
+// The University of British Columbia
+// 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 The University of British Columbia 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 AND CONTRIBUTORS "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 HOLDER 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.
+
+#ifndef _STRINGLIST_H_
+#define _STRINGLIST_H_
+struct stringListPiece
+{
+ const char* stringText;
+
+ stringListPiece* nextString;
+};
+
+
+class stringList
+{
+
+private:
+ //List size
+ int m_size;
+
+ //Start is the first entry, end is the last entry
+ stringListPiece* m_listStart;
+ stringListPiece* m_listEnd;
+public:
+ //constructor
+ stringList();
+
+ //accessors
+ int getSize();
+ stringListPiece* getListStart();
+ stringListPiece* getListEnd();
+
+ //mutator
+ int add(stringListPiece* newString); //add String to list
+ bool remove(int index); //remove string at index
+
+ //print representation
+ void printStringList();
+};
+#endif //_STRINGLIST_H_
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 22:36:56 +0000