diff --git a/examples_ptx/ptxcc/Makefile b/examples_ptx/ptxcc/Makefile new file mode 100644 index 00000000..7db36eec --- /dev/null +++ b/examples_ptx/ptxcc/Makefile @@ -0,0 +1,26 @@ +all: dePTX + +ptxgrammar.cc : ptxgrammar.yy + bison -d -v -t ptxgrammar.yy -o ptxgrammar.cc + +ptx.cc: ptx.ll ptxgrammar.cc + flex -t ptx.ll > ptx.cc + +%.o: %.cc + clang++ -O3 -c $< -o $@ -I/opt/local/include + +%.o: %.cpp + clang++ -O3 -c $< -o $@ -I/opt/local/include + +OBJ= dePTX.o \ + ptx.o \ + ptxgrammar.o + +dePTX: $(OBJ) + clang++ $^ -o $@ -L/opt/local/lib + +clean: + /bin/rm -f dePTX $(OBJ) ptxgrammar.hh ptxgrammar.cc ptx.cc ptxgrammar.output + +$(OBJ): ptxgrammar.cc ptx.cc PTXParser.h PTXLexer.h + diff --git a/examples_ptx/ptxcc/PTXLexer.h b/examples_ptx/ptxcc/PTXLexer.h new file mode 100644 index 00000000..1136fe4d --- /dev/null +++ b/examples_ptx/ptxcc/PTXLexer.h @@ -0,0 +1,40 @@ +#pragma once + +#include +#include + +namespace parser +{ + class PTXLexer; + class PTXParser; +} + +#include "ptxgrammar.hh" + +namespace parser +{ + /*! \brief A wrapper around yyFlexLexer to allow for a local variable */ + class PTXLexer : public ptxFlexLexer + { + public: + YYSTYPE* yylval; + int column; + int nextColumn; + + public: + PTXLexer( std::istream* arg_yyin, + std::ostream* arg_yyout ) : + yyFlexLexer( arg_yyin, arg_yyout ), yylval( 0 ), column( 0 ), + nextColumn( 0 ) { } + + int yylex(); + int yylexPosition() + { + int token = yylex(); + column = nextColumn; + nextColumn = column + strlen( YYText() ); + return token; + } + + }; +} diff --git a/examples_ptx/ptxcc/PTXParser.h b/examples_ptx/ptxcc/PTXParser.h new file mode 100644 index 00000000..f0b6b055 --- /dev/null +++ b/examples_ptx/ptxcc/PTXParser.h @@ -0,0 +1,236 @@ +#pragma once + +#undef yyFlexLexer +#define yyFlexLexer ptxFlexLexer +#include + + +#include "PTXLexer.h" + +#include +#include +#include +namespace ptx +{ + extern int yyparse( parser::PTXLexer&, parser::PTXParser& ); +} + +namespace parser +{ + /*! \brief An implementation of the Parser interface for PTX */ + class PTXParser + { + private: + typedef int token_t; + std::ostream &out; + std::string _identifier; + token_t _dataTypeId; + int _alignment; + + bool isArgumentList, isReturnArgumentList; + struct argument_t + { + token_t type; + std::string name; + int dim; + + argument_t(const token_t _type, const std::string &_name, const int _dim = 1) : + type(_type), name(_name), dim(_dim) {} + }; + std::vector argumentList, returnArgumentList; + std::vector arrayDimensionsList; + + public: + PTXParser(std::ostream &_out) : out(_out) + { + isArgumentList = isReturnArgumentList = false; + _alignment = 1; + } + + void printHeader() + { + std::stringstream s; + s << "template struct __align__(N) b8_t { unsigned char _v[N]; __device__ b8_t() {}; __device__ b8_t (const int value) {}}; \n"; + s << "template struct __align__(2*N) b16_t { unsigned short _v[N]; __device__ b16_t() {}; __device__ b16_t(const int value) {}}; \n"; + s << "struct b8d_t { unsigned char _v[1]; }; \n"; + s << "struct b16d_t { unsigned short _v[1]; }; \n"; + + s << "typedef unsigned int b32_t; \n"; + s << "typedef unsigned int u32_t; \n"; + s << "typedef int s32_t; \n"; + + s << "typedef unsigned long long b64_t; \n"; + s << "typedef unsigned long long u64_t; \n"; + s << "typedef long long s64_t; \n"; + + s << "typedef float f32_t; \n"; + s << "typedef double f64_t; \n"; + s << " \n"; + std::cout << s.str(); + } + +#define LOC YYLTYPE& location + + void identifier(const std::string &s) { _identifier = s; } + void dataTypeId(const token_t token) { _dataTypeId = token; } + void argumentListBegin(LOC) { isArgumentList = true; } + void argumentListEnd (LOC) { isArgumentList = false; } + void returnArgumentListBegin(LOC) { isReturnArgumentList = true; } + void returnArgumentListEnd (LOC) { isReturnArgumentList = false; } + void argumentDeclaration(LOC) + { + assert(arrayDimensionsList.size() <= 1); + const int dim = arrayDimensionsList.empty() ? 1 : arrayDimensionsList[0]; + const argument_t arg(_dataTypeId, _identifier, dim); + if (isArgumentList) + argumentList.push_back(arg); + else if (isReturnArgumentList) + returnArgumentList.push_back(arg); + else + assert(0); + arrayDimensionsList.clear(); + } + void alignment(const int value) { _alignment = value; } + + void arrayDimensions(const int value) + { + arrayDimensionsList.push_back(value); + } + + std::string printArgument(const argument_t arg, const bool printDataType = true) + { + std::stringstream s; + if (printDataType) + s << tokenToDataType(arg.type, arg.dim) << " "; + s << arg.name << " "; + return s.str(); + } + + std::string printArgumentList(const bool printDataType = true) + { + std::stringstream s; + if (argumentList.empty()) return s.str(); + const int n = argumentList.size(); + s << " " << printArgument(argumentList[0], printDataType); + for (int i = 1; i < n; i++) + s << ",\n " << printArgument(argumentList[i], printDataType); + return s.str(); + } + + void visibleEntryDeclaration(const std::string &calleeName, LOC) + { + std::stringstream s; + assert(returnArgumentList.empty()); + s << "extern \"C\" \n"; + s << "__global__ void " << calleeName << " (\n"; + s << printArgumentList(); + s << "\n ) { asm(\" // entry \"); }\n"; + + + /* check if this is an "export" entry */ + const int entryNameLength = calleeName.length(); + const int hostNameLength = std::max(entryNameLength-9,0); + const std::string ___export(&calleeName.c_str()[hostNameLength]); + if (___export.compare("___export") == 0) + { + std::string hostCalleeName; + hostCalleeName.append(calleeName.c_str(), hostNameLength); + s << "/*** host interface ***/\n"; + s << "extern \"C\" \n"; + s << "__host__ void " << hostCalleeName << " (\n"; + s << printArgumentList(); + s << "\n )\n"; + s << "{\n "; + s << calleeName; + s << "<<<1,32>>>(\n"; + s << printArgumentList(false); + s << ");\n"; + s << " cudaDeviceSynchronize(); \n"; + s << "}\n"; + } + s << "\n"; + argumentList.clear(); + + std::cout << s.str(); + } + + void visibleFunctionDeclaration(const std::string &calleeName, LOC) + { + std::stringstream s; + assert(returnArgumentList.size() < 2); + s << "extern \"C\" \n"; + s << "__device__ "; + if (returnArgumentList.empty()) + s << " void "; + else + s << " " << tokenToDataType(returnArgumentList[0].type, returnArgumentList[0].dim); + s << calleeName << " (\n"; + s << printArgumentList(); + + if (returnArgumentList.empty()) + s << "\n ) { asm(\" // function \"); }\n\n"; + else + { + s << "\n ) { asm(\" // function \"); return 0;} /* return value to disable warnings */\n\n"; +// s << "\n ) { asm(\" // function \"); } /* this will generate warrning */\n\n"; + } + + argumentList.clear(); + returnArgumentList.clear(); + + std::cout << s.str(); + } + + void visibleInitializableDeclaration(const std::string &name, LOC) + { + assert(arrayDimensionsList.size() == 1); + std::stringstream s; + s << "extern \"C\" __device__ "; + if (_alignment > 0) + s << "__attribute__((aligned(" << _alignment << "))) "; + s << tokenToDataType(_dataTypeId, 0); + s << name << "[" << arrayDimensionsList[0] << "] = {0};\n\n"; + std::cout << s.str(); + arrayDimensionsList.clear(); + } + +#undef LOC + + std::string tokenToDataType( token_t token , int dim) + { + std::stringstream s; + switch( token ) + { + case TOKEN_B8: + if (dim > 0) s << "b8_t<"< "; + else s << "b8d_t "; + break; + case TOKEN_U8: assert(0); s << "u8_t "; break; + case TOKEN_S8: assert(0); s << "s8_t "; break; + + case TOKEN_B16: + if (dim > 0) s << "b16_t<"< "; + else s << "b16d_t "; + break; + case TOKEN_U16: assert(0); s << "u16_t "; break; + case TOKEN_S16: assert(0); s << "s16_t "; break; + + case TOKEN_B32: assert(dim == 1); s << "b32_t "; break; + case TOKEN_U32: assert(dim == 1); s << "u32_t "; break; + case TOKEN_S32: assert(dim == 1); s << "s32_t "; break; + + case TOKEN_B64: assert(dim == 1); s << "b64_t "; break; + case TOKEN_U64: assert(dim == 1); s << "u64_t "; break; + case TOKEN_S64: assert(dim == 1); s << "s64_t "; break; + + case TOKEN_F32: assert(dim == 1); s << "f32_t "; break; + case TOKEN_F64: assert(dim == 1); s << "f64_t "; break; + default: std::cerr << "token= " << token<< std::endl; assert(0); + } + + return s.str(); + } + }; +} + + diff --git a/examples_ptx/ptxcc/dePTX.cpp b/examples_ptx/ptxcc/dePTX.cpp new file mode 100644 index 00000000..5d9c5c98 --- /dev/null +++ b/examples_ptx/ptxcc/dePTX.cpp @@ -0,0 +1,24 @@ +#include +#include +#include +#include +#include "PTXParser.h" + +int main(int argc, char * argv[]) +{ + // open a file handle to a particular file: + std::istream & input = std::cin; + std::ostream & error = std::cerr; + std::ostream & output = std::cout; + parser::PTXLexer lexer(&input, &error); + parser::PTXParser state(output); + + // parse through the input until there is no more: + // + + do { + ptx::yyparse(lexer, state); + } + while (!input.eof()); + +} diff --git a/examples_ptx/ptxcc/ptx.ll b/examples_ptx/ptxcc/ptx.ll new file mode 100644 index 00000000..04d546fa --- /dev/null +++ b/examples_ptx/ptxcc/ptx.ll @@ -0,0 +1,76 @@ +%option yylineno +%option noyywrap +%option yyclass="parser::PTXLexer" +%option prefix="ptx" +%option c++ + +%{ +#include "PTXLexer.h" +#include +#include +#include +#ifdef LLSETTOKEN +#error "TOKEN is defined" +#endif +#define LLSETTOKEN(tok) yylval->ivalue = tok; return tok; +%} + +COMMENT ("//"[^\n]*) +TAB [\t]* + +%% +{COMMENT} {nextColumn += strlen(yytext); /* lCppComment(&yylloc); */ } +".version" { return TOKEN_VERSION; } +".target" { return TOKEN_TARGET; } +".address_size" { return TOKEN_ADDRESS_SIZE; } +".func" { return TOKEN_FUNC; } +".entry" { return TOKEN_ENTRY; } +".align" { return TOKEN_ALIGN; } +".visible" { return TOKEN_VISIBLE; } +".global" { return TOKEN_GLOBAL; } +".param" { return TOKEN_PARAM; } +".b8" { LLSETTOKEN( TOKEN_B8);} +".b16" { LLSETTOKEN( TOKEN_B16);} +".b32" { LLSETTOKEN( TOKEN_B32);} +".b64" { LLSETTOKEN( TOKEN_B64);} +".u8" { LLSETTOKEN( TOKEN_U8);} +".u16" { LLSETTOKEN( TOKEN_U16);} +".u32" { LLSETTOKEN( TOKEN_U32);} +".u64" { LLSETTOKEN( TOKEN_U64);} +".s8" { LLSETTOKEN( TOKEN_S8);} +".s16" { LLSETTOKEN( TOKEN_S16);} +".s32" { LLSETTOKEN( TOKEN_S32);} +".s64" { LLSETTOKEN( TOKEN_S64);} +".f32" { LLSETTOKEN( TOKEN_F32);} +".f64" { LLSETTOKEN( TOKEN_F64);} +"[" { return '[';} +"]" { return ']';} +"(" { return '(';} +")" { return ')';} +"," { return ',';} +[0-9]+\.[0-9]+ { yylval->fvalue = atof(yytext); return TOKEN_FLOAT; } +[0-9]+ { yylval->ivalue = atoi(yytext); return TOKEN_INT; } +[a-zA-Z0-9_]+ { strcpy(yylval->svalue, yytext); return TOKEN_STRING;} +\n { + // yylloc.last_line++; +// yylloc.last_column = 1; + nextColumn = 1; +} +. ; +%% + +/** Handle a C++-style comment--eat everything up until the end of the line. + */ +#if 0 +static void +lCppComment(SourcePos *pos) { + char c; + do { + c = yyinput(); + } while (c != 0 && c != '\n'); + if (c == '\n') { + pos->last_line++; + pos->last_column = 1; + } +} +#endif diff --git a/examples_ptx/ptxcc/ptxcc b/examples_ptx/ptxcc/ptxcc new file mode 100755 index 00000000..0f7e384e --- /dev/null +++ b/examples_ptx/ptxcc/ptxcc @@ -0,0 +1,18 @@ +#!/bin/sh + +PTXSRC=$1 +PTXCU=$1___tmp_ptx.cu +PTXSH=$1___tmp_ptx.sh + +NVCCPARM=${@:2} + +DEPTX=dePTX +NVCC=nvcc + +$DEPTX < $PTXSRC > $PTXCU && +$NVCC -arch=sm_35 -dc $NVCCPARM -dryrun $PTXCU 2>&1 | \ + sed 's/\#\$//g'| \ + awk '{ if ($1 == "LIBRARIES=") print $1$2; else if ($1 == "cicc") print "cp '$PTXSRC'", $NF; else print $0 }' > $PTXSH && +sh $PTXSH + +# rm $PTXCU $PTXSH diff --git a/examples_ptx/ptxcc/ptxgrammar.cpp b/examples_ptx/ptxcc/ptxgrammar.cpp new file mode 100644 index 00000000..2419e654 --- /dev/null +++ b/examples_ptx/ptxcc/ptxgrammar.cpp @@ -0,0 +1,1929 @@ + +/* A Bison parser, made by GNU Bison 2.4.1. */ + +/* Skeleton implementation for Bison's Yacc-like parsers in C + + Copyright (C) 1984, 1989, 1990, 2000, 2001, 2002, 2003, 2004, 2005, 2006 + Free Software Foundation, Inc. + + This program is free software: you can redistribute it and/or modify + it under the terms of the GNU General Public License as published by + the Free Software Foundation, either version 3 of the License, or + (at your option) any later version. + + This program is distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + GNU General Public License for more details. + + You should have received a copy of the GNU General Public License + along with this program. If not, see . */ + +/* As a special exception, you may create a larger work that contains + part or all of the Bison parser skeleton and distribute that work + under terms of your choice, so long as that work isn't itself a + parser generator using the skeleton or a modified version thereof + as a parser skeleton. Alternatively, if you modify or redistribute + the parser skeleton itself, you may (at your option) remove this + special exception, which will cause the skeleton and the resulting + Bison output files to be licensed under the GNU General Public + License without this special exception. + + This special exception was added by the Free Software Foundation in + version 2.2 of Bison. */ + +/* C LALR(1) parser skeleton written by Richard Stallman, by + simplifying the original so-called "semantic" parser. */ + +/* All symbols defined below should begin with yy or YY, to avoid + infringing on user name space. This should be done even for local + variables, as they might otherwise be expanded by user macros. + There are some unavoidable exceptions within include files to + define necessary library symbols; they are noted "INFRINGES ON + USER NAME SPACE" below. */ + +/* Identify Bison output. */ +#define YYBISON 1 + +/* Bison version. */ +#define YYBISON_VERSION "2.4.1" + +/* Skeleton name. */ +#define YYSKELETON_NAME "yacc.c" + +/* Pure parsers. */ +#define YYPURE 1 + +/* Push parsers. */ +#define YYPUSH 0 + +/* Pull parsers. */ +#define YYPULL 1 + +/* Using locations. */ +#define YYLSP_NEEDED 1 + + + +/* Copy the first part of user declarations. */ + +/* Line 189 of yacc.c */ +#line 3 "ptxgrammar.yy" + + #include + #include "PTXParser.h" + #include "PTXLexer.h" + #include + #include + #include + + #define YYERROR_VERBOSE 1 + + #ifdef REPORT_BASE + #undef REPORT_BASE + #endif + + #define REPORT_BASE 0 + + namespace ptx + { + + int yylex( YYSTYPE* token, YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state ); + void yyerror( YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state, char const* message ); + + std::string yyTypeToString( int ); + + + +/* Line 189 of yacc.c */ +#line 102 "ptxgrammar.cc" + +/* Enabling traces. */ +#ifndef YYDEBUG +# define YYDEBUG 1 +#endif + +/* Enabling verbose error messages. */ +#ifdef YYERROR_VERBOSE +# undef YYERROR_VERBOSE +# define YYERROR_VERBOSE 1 +#else +# define YYERROR_VERBOSE 0 +#endif + +/* Enabling the token table. */ +#ifndef YYTOKEN_TABLE +# define YYTOKEN_TABLE 0 +#endif + + +/* Tokens. */ +#ifndef YYTOKENTYPE +# define YYTOKENTYPE + /* Put the tokens into the symbol table, so that GDB and other debuggers + know about them. */ + enum yytokentype { + TOKEN_VERSION = 258, + TOKEN_TARGET = 259, + TOKEN_ADDRESS_SIZE = 260, + TOKEN_VISIBLE = 261, + TOKEN_FUNC = 262, + TOKEN_ENTRY = 263, + TOKEN_PARAM = 264, + TOKEN_ALIGN = 265, + TOKEN_GLOBAL = 266, + TOKEN_B8 = 267, + TOKEN_B16 = 268, + TOKEN_B32 = 269, + TOKEN_B64 = 270, + TOKEN_U8 = 271, + TOKEN_U16 = 272, + TOKEN_U32 = 273, + TOKEN_U64 = 274, + TOKEN_S8 = 275, + TOKEN_S16 = 276, + TOKEN_S32 = 277, + TOKEN_S64 = 278, + TOKEN_F32 = 279, + TOKEN_F64 = 280, + TOKEN_INT = 281, + TOKEN_FLOAT = 282, + TOKEN_STRING = 283 + }; +#endif + + + +#if ! defined YYSTYPE && ! defined YYSTYPE_IS_DECLARED +typedef union YYSTYPE +{ + +/* Line 214 of yacc.c */ +#line 32 "ptxgrammar.yy" + + char svalue[1024]; + double fvalue; + int ivalue; + unsigned int uvalue; + + + +/* Line 214 of yacc.c */ +#line 175 "ptxgrammar.cc" +} YYSTYPE; +# define YYSTYPE_IS_TRIVIAL 1 +# define yystype YYSTYPE /* obsolescent; will be withdrawn */ +# define YYSTYPE_IS_DECLARED 1 +#endif + +#if ! defined YYLTYPE && ! defined YYLTYPE_IS_DECLARED +typedef struct YYLTYPE +{ + int first_line; + int first_column; + int last_line; + int last_column; +} YYLTYPE; +# define yyltype YYLTYPE /* obsolescent; will be withdrawn */ +# define YYLTYPE_IS_DECLARED 1 +# define YYLTYPE_IS_TRIVIAL 1 +#endif + + +/* Copy the second part of user declarations. */ + + +/* Line 264 of yacc.c */ +#line 200 "ptxgrammar.cc" + +#ifdef short +# undef short +#endif + +#ifdef YYTYPE_UINT8 +typedef YYTYPE_UINT8 yytype_uint8; +#else +typedef unsigned char yytype_uint8; +#endif + +#ifdef YYTYPE_INT8 +typedef YYTYPE_INT8 yytype_int8; +#elif (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +typedef signed char yytype_int8; +#else +typedef short int yytype_int8; +#endif + +#ifdef YYTYPE_UINT16 +typedef YYTYPE_UINT16 yytype_uint16; +#else +typedef unsigned short int yytype_uint16; +#endif + +#ifdef YYTYPE_INT16 +typedef YYTYPE_INT16 yytype_int16; +#else +typedef short int yytype_int16; +#endif + +#ifndef YYSIZE_T +# ifdef __SIZE_TYPE__ +# define YYSIZE_T __SIZE_TYPE__ +# elif defined size_t +# define YYSIZE_T size_t +# elif ! defined YYSIZE_T && (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +# include /* INFRINGES ON USER NAME SPACE */ +# define YYSIZE_T size_t +# else +# define YYSIZE_T unsigned int +# endif +#endif + +#define YYSIZE_MAXIMUM ((YYSIZE_T) -1) + +#ifndef YY_ +# if YYENABLE_NLS +# if ENABLE_NLS +# include /* INFRINGES ON USER NAME SPACE */ +# define YY_(msgid) dgettext ("bison-runtime", msgid) +# endif +# endif +# ifndef YY_ +# define YY_(msgid) msgid +# endif +#endif + +/* Suppress unused-variable warnings by "using" E. */ +#if ! defined lint || defined __GNUC__ +# define YYUSE(e) ((void) (e)) +#else +# define YYUSE(e) /* empty */ +#endif + +/* Identity function, used to suppress warnings about constant conditions. */ +#ifndef lint +# define YYID(n) (n) +#else +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static int +YYID (int yyi) +#else +static int +YYID (yyi) + int yyi; +#endif +{ + return yyi; +} +#endif + +#if ! defined yyoverflow || YYERROR_VERBOSE + +/* The parser invokes alloca or malloc; define the necessary symbols. */ + +# ifdef YYSTACK_USE_ALLOCA +# if YYSTACK_USE_ALLOCA +# ifdef __GNUC__ +# define YYSTACK_ALLOC __builtin_alloca +# elif defined __BUILTIN_VA_ARG_INCR +# include /* INFRINGES ON USER NAME SPACE */ +# elif defined _AIX +# define YYSTACK_ALLOC __alloca +# elif defined _MSC_VER +# include /* INFRINGES ON USER NAME SPACE */ +# define alloca _alloca +# else +# define YYSTACK_ALLOC alloca +# if ! defined _ALLOCA_H && ! defined _STDLIB_H && (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +# include /* INFRINGES ON USER NAME SPACE */ +# ifndef _STDLIB_H +# define _STDLIB_H 1 +# endif +# endif +# endif +# endif +# endif + +# ifdef YYSTACK_ALLOC + /* Pacify GCC's `empty if-body' warning. */ +# define YYSTACK_FREE(Ptr) do { /* empty */; } while (YYID (0)) +# ifndef YYSTACK_ALLOC_MAXIMUM + /* The OS might guarantee only one guard page at the bottom of the stack, + and a page size can be as small as 4096 bytes. So we cannot safely + invoke alloca (N) if N exceeds 4096. Use a slightly smaller number + to allow for a few compiler-allocated temporary stack slots. */ +# define YYSTACK_ALLOC_MAXIMUM 4032 /* reasonable circa 2006 */ +# endif +# else +# define YYSTACK_ALLOC YYMALLOC +# define YYSTACK_FREE YYFREE +# ifndef YYSTACK_ALLOC_MAXIMUM +# define YYSTACK_ALLOC_MAXIMUM YYSIZE_MAXIMUM +# endif +# if (defined __cplusplus && ! defined _STDLIB_H \ + && ! ((defined YYMALLOC || defined malloc) \ + && (defined YYFREE || defined free))) +# include /* INFRINGES ON USER NAME SPACE */ +# ifndef _STDLIB_H +# define _STDLIB_H 1 +# endif +# endif +# ifndef YYMALLOC +# define YYMALLOC malloc +# if ! defined malloc && ! defined _STDLIB_H && (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +void *malloc (YYSIZE_T); /* INFRINGES ON USER NAME SPACE */ +# endif +# endif +# ifndef YYFREE +# define YYFREE free +# if ! defined free && ! defined _STDLIB_H && (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +void free (void *); /* INFRINGES ON USER NAME SPACE */ +# endif +# endif +# endif +#endif /* ! defined yyoverflow || YYERROR_VERBOSE */ + + +#if (! defined yyoverflow \ + && (! defined __cplusplus \ + || (defined YYLTYPE_IS_TRIVIAL && YYLTYPE_IS_TRIVIAL \ + && defined YYSTYPE_IS_TRIVIAL && YYSTYPE_IS_TRIVIAL))) + +/* A type that is properly aligned for any stack member. */ +union yyalloc +{ + yytype_int16 yyss_alloc; + YYSTYPE yyvs_alloc; + YYLTYPE yyls_alloc; +}; + +/* The size of the maximum gap between one aligned stack and the next. */ +# define YYSTACK_GAP_MAXIMUM (sizeof (union yyalloc) - 1) + +/* The size of an array large to enough to hold all stacks, each with + N elements. */ +# define YYSTACK_BYTES(N) \ + ((N) * (sizeof (yytype_int16) + sizeof (YYSTYPE) + sizeof (YYLTYPE)) \ + + 2 * YYSTACK_GAP_MAXIMUM) + +/* Copy COUNT objects from FROM to TO. The source and destination do + not overlap. */ +# ifndef YYCOPY +# if defined __GNUC__ && 1 < __GNUC__ +# define YYCOPY(To, From, Count) \ + __builtin_memcpy (To, From, (Count) * sizeof (*(From))) +# else +# define YYCOPY(To, From, Count) \ + do \ + { \ + YYSIZE_T yyi; \ + for (yyi = 0; yyi < (Count); yyi++) \ + (To)[yyi] = (From)[yyi]; \ + } \ + while (YYID (0)) +# endif +# endif + +/* Relocate STACK from its old location to the new one. The + local variables YYSIZE and YYSTACKSIZE give the old and new number of + elements in the stack, and YYPTR gives the new location of the + stack. Advance YYPTR to a properly aligned location for the next + stack. */ +# define YYSTACK_RELOCATE(Stack_alloc, Stack) \ + do \ + { \ + YYSIZE_T yynewbytes; \ + YYCOPY (&yyptr->Stack_alloc, Stack, yysize); \ + Stack = &yyptr->Stack_alloc; \ + yynewbytes = yystacksize * sizeof (*Stack) + YYSTACK_GAP_MAXIMUM; \ + yyptr += yynewbytes / sizeof (*yyptr); \ + } \ + while (YYID (0)) + +#endif + +/* YYFINAL -- State number of the termination state. */ +#define YYFINAL 6 +/* YYLAST -- Last index in YYTABLE. */ +#define YYLAST 97 + +/* YYNTOKENS -- Number of terminals. */ +#define YYNTOKENS 34 +/* YYNNTS -- Number of nonterminals. */ +#define YYNNTS 28 +/* YYNRULES -- Number of rules. */ +#define YYNRULES 66 +/* YYNRULES -- Number of states. */ +#define YYNSTATES 92 + +/* YYTRANSLATE(YYLEX) -- Bison symbol number corresponding to YYLEX. */ +#define YYUNDEFTOK 2 +#define YYMAXUTOK 283 + +#define YYTRANSLATE(YYX) \ + ((unsigned int) (YYX) <= YYMAXUTOK ? yytranslate[YYX] : YYUNDEFTOK) + +/* YYTRANSLATE[YYLEX] -- Bison symbol number corresponding to YYLEX. */ +static const yytype_uint8 yytranslate[] = +{ + 0, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 31, 32, 2, 2, 33, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 29, 2, 30, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, + 2, 2, 2, 2, 2, 2, 1, 2, 3, 4, + 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, + 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, + 25, 26, 27, 28 +}; + +#if YYDEBUG +/* YYPRHS[YYN] -- Index of the first RHS symbol of rule number YYN in + YYRHS. */ +static const yytype_uint8 yyprhs[] = +{ + 0, 0, 3, 6, 10, 13, 16, 19, 21, 23, + 25, 27, 29, 31, 33, 35, 37, 39, 41, 43, + 45, 47, 49, 51, 53, 55, 57, 59, 61, 63, + 65, 67, 69, 71, 73, 75, 77, 80, 82, 85, + 87, 90, 92, 95, 97, 101, 102, 104, 106, 108, + 111, 113, 116, 121, 123, 125, 127, 128, 132, 136, + 141, 143, 145, 149, 151, 152, 158 +}; + +/* YYRHS -- A `-1'-separated list of the rules' RHS. */ +static const yytype_int8 yyrhs[] = +{ + 35, 0, -1, 36, 43, -1, 37, 38, 39, -1, + 3, 27, -1, 4, 28, -1, 5, 26, -1, 16, + -1, 17, -1, 18, -1, 19, -1, 20, -1, 21, + -1, 22, -1, 23, -1, 12, -1, 13, -1, 14, + -1, 15, -1, 24, -1, 25, -1, 40, -1, 10, + -1, 9, -1, 40, -1, 28, -1, 27, -1, 26, + -1, 7, -1, 8, -1, 11, -1, 29, -1, 30, + -1, 31, -1, 32, -1, 33, -1, 43, 60, -1, + 60, -1, 43, 55, -1, 55, -1, 43, 61, -1, + 61, -1, 43, 42, -1, 42, -1, 29, 26, 30, + -1, -1, 44, -1, 28, -1, 9, -1, 10, 26, + -1, 41, -1, 48, 41, -1, 47, 49, 46, 45, + -1, 31, -1, 32, -1, 50, -1, -1, 53, 33, + 50, -1, 51, 53, 52, -1, 6, 8, 46, 54, + -1, 31, -1, 32, -1, 56, 53, 57, -1, 58, + -1, -1, 6, 7, 59, 46, 54, -1, 6, 11, + 49, 46, 44, -1 +}; + +/* YYRLINE[YYN] -- source line where rule number YYN was defined. */ +static const yytype_uint8 yyrline[] = +{ + 0, 71, 71, 74, 81, 83, 85, 89, 89, 89, + 89, 90, 90, 90, 90, 91, 91, 91, 91, 92, + 92, 94, 97, 98, 99, 100, 100, 100, 101, 101, + 102, 103, 104, 105, 106, 107, 110, 110, 111, 111, + 112, 112, 113, 113, 117, 120, 121, 123, 124, 126, + 127, 128, 130, 136, 137, 138, 139, 140, 141, 143, + 148, 149, 150, 151, 151, 152, 158 +}; +#endif + +#if YYDEBUG || YYERROR_VERBOSE || YYTOKEN_TABLE +/* YYTNAME[SYMBOL-NUM] -- String name of the symbol SYMBOL-NUM. + First, the terminals, then, starting at YYNTOKENS, nonterminals. */ +static const char *const yytname[] = +{ + "$end", "error", "$undefined", "TOKEN_VERSION", "TOKEN_TARGET", + "TOKEN_ADDRESS_SIZE", "TOKEN_VISIBLE", "TOKEN_FUNC", "TOKEN_ENTRY", + "TOKEN_PARAM", "TOKEN_ALIGN", "TOKEN_GLOBAL", "TOKEN_B8", "TOKEN_B16", + "TOKEN_B32", "TOKEN_B64", "TOKEN_U8", "TOKEN_U16", "TOKEN_U32", + "TOKEN_U64", "TOKEN_S8", "TOKEN_S16", "TOKEN_S32", "TOKEN_S64", + "TOKEN_F32", "TOKEN_F64", "TOKEN_INT", "TOKEN_FLOAT", "TOKEN_STRING", + "'['", "']'", "'('", "')'", "','", "$accept", "ptxsource", "header", + "version", "target", "address_size", "dataTypeId", "dataType", + "anytoken", "ptxbody", "arrayDimensionSet", "arrayDimensions", + "identifier", "parameter", "alignment", "addressableVariablePrefix", + "argumentDeclaration", "argumentListBegin", "argumentListEnd", + "argumentListBody", "argumentList", "visibleEntryDeclaration", + "returnArgumentListBegin", "returnArgumentListEnd", "returnArgumentList", + "optionalReturnArgumentList", "visibleFunctionDeclaration", + "visibleInitializableDeclaration", 0 +}; +#endif + +# ifdef YYPRINT +/* YYTOKNUM[YYLEX-NUM] -- Internal token number corresponding to + token YYLEX-NUM. */ +static const yytype_uint16 yytoknum[] = +{ + 0, 256, 257, 258, 259, 260, 261, 262, 263, 264, + 265, 266, 267, 268, 269, 270, 271, 272, 273, 274, + 275, 276, 277, 278, 279, 280, 281, 282, 283, 91, + 93, 40, 41, 44 +}; +# endif + +/* YYR1[YYN] -- Symbol number of symbol that rule YYN derives. */ +static const yytype_uint8 yyr1[] = +{ + 0, 34, 35, 36, 37, 38, 39, 40, 40, 40, + 40, 40, 40, 40, 40, 40, 40, 40, 40, 40, + 40, 41, 42, 42, 42, 42, 42, 42, 42, 42, + 42, 42, 42, 42, 42, 42, 43, 43, 43, 43, + 43, 43, 43, 43, 44, 45, 45, 46, 47, 48, + 49, 49, 50, 51, 52, 53, 53, 53, 54, 55, + 56, 57, 58, 59, 59, 60, 61 +}; + +/* YYR2[YYN] -- Number of symbols composing right hand side of rule YYN. */ +static const yytype_uint8 yyr2[] = +{ + 0, 2, 2, 3, 2, 2, 2, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, + 1, 1, 1, 1, 1, 1, 2, 1, 2, 1, + 2, 1, 2, 1, 3, 0, 1, 1, 1, 2, + 1, 2, 4, 1, 1, 1, 0, 3, 3, 4, + 1, 1, 3, 1, 0, 5, 5 +}; + +/* YYDEFACT[STATE-NAME] -- Default rule to reduce with in state + STATE-NUM when YYTABLE doesn't specify something else to do. Zero + means the default is an error. */ +static const yytype_uint8 yydefact[] = +{ + 0, 0, 0, 0, 0, 4, 1, 0, 28, 29, + 23, 22, 30, 15, 16, 17, 18, 7, 8, 9, + 10, 11, 12, 13, 14, 19, 20, 27, 26, 25, + 31, 32, 33, 34, 35, 24, 43, 2, 39, 37, + 41, 0, 0, 64, 0, 0, 42, 38, 36, 40, + 5, 0, 3, 60, 56, 63, 0, 47, 0, 0, + 21, 50, 0, 0, 6, 48, 0, 55, 0, 0, + 53, 56, 59, 49, 51, 0, 0, 61, 0, 62, + 65, 0, 0, 66, 45, 57, 54, 58, 0, 46, + 52, 44 +}; + +/* YYDEFGOTO[NTERM-NUM]. */ +static const yytype_int8 yydefgoto[] = +{ + -1, 2, 3, 4, 42, 52, 60, 61, 36, 37, + 83, 90, 58, 66, 62, 63, 67, 71, 87, 68, + 72, 38, 54, 79, 55, 56, 39, 40 +}; + +/* YYPACT[STATE-NUM] -- Index in YYTABLE of the portion describing + STATE-NUM. */ +#define YYPACT_NINF -28 +static const yytype_int8 yypact[] = +{ + 27, 6, 41, -5, 39, -28, -28, 24, -28, -28, + -28, -28, -28, -28, -28, -28, -28, -28, -28, -28, + -28, -28, -28, -28, -28, -28, -28, -28, -28, -28, + -28, -28, -28, -28, -28, -28, -28, -5, -28, -28, + -28, 14, 40, 13, 18, 38, -28, -28, -28, -28, + -28, 21, -28, -28, 69, -28, 18, -28, 48, 54, + -28, -28, 52, 18, -28, -28, 38, -28, 5, 48, + -28, 69, -28, -28, -28, 53, 18, -28, 69, -28, + -28, 7, 55, -28, 53, -28, -28, -28, 56, -28, + -28, -28 +}; + +/* YYPGOTO[NTERM-NUM]. */ +static const yytype_int8 yypgoto[] = +{ + -28, -28, -28, -28, -28, -28, -3, 22, 46, -28, + 1, -28, -27, -28, -28, 23, 9, -28, -28, 17, + 25, 58, -28, -28, -28, -28, 59, 60 +}; + +/* YYTABLE[YYPACT[STATE-NUM]]. What to do in state STATE-NUM. If + positive, shift that token. If negative, reduce the rule which + number is the opposite. If zero, do what YYDEFACT says. + If YYTABLE_NINF, syntax error. */ +#define YYTABLE_NINF -1 +static const yytype_uint8 yytable[] = +{ + 35, 7, 8, 9, 10, 11, 12, 13, 14, 15, + 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, + 26, 27, 28, 29, 30, 31, 32, 33, 34, 69, + 1, 43, 44, 5, 35, 45, 75, 77, 78, 86, + 78, 6, 50, 41, 53, 51, 57, 64, 59, 84, + 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, + 23, 24, 25, 26, 13, 14, 15, 16, 17, 18, + 19, 20, 21, 22, 23, 24, 25, 26, 65, 70, + 73, 88, 82, 46, 74, 89, 91, 85, 81, 76, + 0, 0, 0, 0, 80, 47, 48, 49 +}; + +static const yytype_int8 yycheck[] = +{ + 3, 6, 7, 8, 9, 10, 11, 12, 13, 14, + 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, + 25, 26, 27, 28, 29, 30, 31, 32, 33, 56, + 3, 7, 8, 27, 37, 11, 63, 32, 33, 32, + 33, 0, 28, 4, 31, 5, 28, 26, 10, 76, + 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, + 22, 23, 24, 25, 12, 13, 14, 15, 16, 17, + 18, 19, 20, 21, 22, 23, 24, 25, 9, 31, + 26, 26, 29, 37, 62, 84, 30, 78, 71, 66, + -1, -1, -1, -1, 69, 37, 37, 37 +}; + +/* YYSTOS[STATE-NUM] -- The (internal number of the) accessing + symbol of state STATE-NUM. */ +static const yytype_uint8 yystos[] = +{ + 0, 3, 35, 36, 37, 27, 0, 6, 7, 8, + 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, + 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, + 29, 30, 31, 32, 33, 40, 42, 43, 55, 60, + 61, 4, 38, 7, 8, 11, 42, 55, 60, 61, + 28, 5, 39, 31, 56, 58, 59, 28, 46, 10, + 40, 41, 48, 49, 26, 9, 47, 50, 53, 46, + 31, 51, 54, 26, 41, 46, 49, 32, 33, 57, + 54, 53, 29, 44, 46, 50, 32, 52, 26, 44, + 45, 30 +}; + +#define yyerrok (yyerrstatus = 0) +#define yyclearin (yychar = YYEMPTY) +#define YYEMPTY (-2) +#define YYEOF 0 + +#define YYACCEPT goto yyacceptlab +#define YYABORT goto yyabortlab +#define YYERROR goto yyerrorlab + + +/* Like YYERROR except do call yyerror. This remains here temporarily + to ease the transition to the new meaning of YYERROR, for GCC. + Once GCC version 2 has supplanted version 1, this can go. */ + +#define YYFAIL goto yyerrlab + +#define YYRECOVERING() (!!yyerrstatus) + +#define YYBACKUP(Token, Value) \ +do \ + if (yychar == YYEMPTY && yylen == 1) \ + { \ + yychar = (Token); \ + yylval = (Value); \ + yytoken = YYTRANSLATE (yychar); \ + YYPOPSTACK (1); \ + goto yybackup; \ + } \ + else \ + { \ + yyerror (&yylloc, lexer, state, YY_("syntax error: cannot back up")); \ + YYERROR; \ + } \ +while (YYID (0)) + + +#define YYTERROR 1 +#define YYERRCODE 256 + + +/* YYLLOC_DEFAULT -- Set CURRENT to span from RHS[1] to RHS[N]. + If N is 0, then set CURRENT to the empty location which ends + the previous symbol: RHS[0] (always defined). */ + +#define YYRHSLOC(Rhs, K) ((Rhs)[K]) +#ifndef YYLLOC_DEFAULT +# define YYLLOC_DEFAULT(Current, Rhs, N) \ + do \ + if (YYID (N)) \ + { \ + (Current).first_line = YYRHSLOC (Rhs, 1).first_line; \ + (Current).first_column = YYRHSLOC (Rhs, 1).first_column; \ + (Current).last_line = YYRHSLOC (Rhs, N).last_line; \ + (Current).last_column = YYRHSLOC (Rhs, N).last_column; \ + } \ + else \ + { \ + (Current).first_line = (Current).last_line = \ + YYRHSLOC (Rhs, 0).last_line; \ + (Current).first_column = (Current).last_column = \ + YYRHSLOC (Rhs, 0).last_column; \ + } \ + while (YYID (0)) +#endif + + +/* YY_LOCATION_PRINT -- Print the location on the stream. + This macro was not mandated originally: define only if we know + we won't break user code: when these are the locations we know. */ + +#ifndef YY_LOCATION_PRINT +# if YYLTYPE_IS_TRIVIAL +# define YY_LOCATION_PRINT(File, Loc) \ + fprintf (File, "%d.%d-%d.%d", \ + (Loc).first_line, (Loc).first_column, \ + (Loc).last_line, (Loc).last_column) +# else +# define YY_LOCATION_PRINT(File, Loc) ((void) 0) +# endif +#endif + + +/* YYLEX -- calling `yylex' with the right arguments. */ + +#ifdef YYLEX_PARAM +# define YYLEX yylex (&yylval, &yylloc, YYLEX_PARAM) +#else +# define YYLEX yylex (&yylval, &yylloc, lexer, state) +#endif + +/* Enable debugging if requested. */ +#if YYDEBUG + +# ifndef YYFPRINTF +# include /* INFRINGES ON USER NAME SPACE */ +# define YYFPRINTF fprintf +# endif + +# define YYDPRINTF(Args) \ +do { \ + if (yydebug) \ + YYFPRINTF Args; \ +} while (YYID (0)) + +# define YY_SYMBOL_PRINT(Title, Type, Value, Location) \ +do { \ + if (yydebug) \ + { \ + YYFPRINTF (stderr, "%s ", Title); \ + yy_symbol_print (stderr, \ + Type, Value, Location, lexer, state); \ + YYFPRINTF (stderr, "\n"); \ + } \ +} while (YYID (0)) + + +/*--------------------------------. +| Print this symbol on YYOUTPUT. | +`--------------------------------*/ + +/*ARGSUSED*/ +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static void +yy_symbol_value_print (FILE *yyoutput, int yytype, YYSTYPE const * const yyvaluep, YYLTYPE const * const yylocationp, parser::PTXLexer& lexer, parser::PTXParser& state) +#else +static void +yy_symbol_value_print (yyoutput, yytype, yyvaluep, yylocationp, lexer, state) + FILE *yyoutput; + int yytype; + YYSTYPE const * const yyvaluep; + YYLTYPE const * const yylocationp; + parser::PTXLexer& lexer; + parser::PTXParser& state; +#endif +{ + if (!yyvaluep) + return; + YYUSE (yylocationp); + YYUSE (lexer); + YYUSE (state); +# ifdef YYPRINT + if (yytype < YYNTOKENS) + YYPRINT (yyoutput, yytoknum[yytype], *yyvaluep); +# else + YYUSE (yyoutput); +# endif + switch (yytype) + { + default: + break; + } +} + + +/*--------------------------------. +| Print this symbol on YYOUTPUT. | +`--------------------------------*/ + +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static void +yy_symbol_print (FILE *yyoutput, int yytype, YYSTYPE const * const yyvaluep, YYLTYPE const * const yylocationp, parser::PTXLexer& lexer, parser::PTXParser& state) +#else +static void +yy_symbol_print (yyoutput, yytype, yyvaluep, yylocationp, lexer, state) + FILE *yyoutput; + int yytype; + YYSTYPE const * const yyvaluep; + YYLTYPE const * const yylocationp; + parser::PTXLexer& lexer; + parser::PTXParser& state; +#endif +{ + if (yytype < YYNTOKENS) + YYFPRINTF (yyoutput, "token %s (", yytname[yytype]); + else + YYFPRINTF (yyoutput, "nterm %s (", yytname[yytype]); + + YY_LOCATION_PRINT (yyoutput, *yylocationp); + YYFPRINTF (yyoutput, ": "); + yy_symbol_value_print (yyoutput, yytype, yyvaluep, yylocationp, lexer, state); + YYFPRINTF (yyoutput, ")"); +} + +/*------------------------------------------------------------------. +| yy_stack_print -- Print the state stack from its BOTTOM up to its | +| TOP (included). | +`------------------------------------------------------------------*/ + +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static void +yy_stack_print (yytype_int16 *yybottom, yytype_int16 *yytop) +#else +static void +yy_stack_print (yybottom, yytop) + yytype_int16 *yybottom; + yytype_int16 *yytop; +#endif +{ + YYFPRINTF (stderr, "Stack now"); + for (; yybottom <= yytop; yybottom++) + { + int yybot = *yybottom; + YYFPRINTF (stderr, " %d", yybot); + } + YYFPRINTF (stderr, "\n"); +} + +# define YY_STACK_PRINT(Bottom, Top) \ +do { \ + if (yydebug) \ + yy_stack_print ((Bottom), (Top)); \ +} while (YYID (0)) + + +/*------------------------------------------------. +| Report that the YYRULE is going to be reduced. | +`------------------------------------------------*/ + +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static void +yy_reduce_print (YYSTYPE *yyvsp, YYLTYPE *yylsp, int yyrule, parser::PTXLexer& lexer, parser::PTXParser& state) +#else +static void +yy_reduce_print (yyvsp, yylsp, yyrule, lexer, state) + YYSTYPE *yyvsp; + YYLTYPE *yylsp; + int yyrule; + parser::PTXLexer& lexer; + parser::PTXParser& state; +#endif +{ + int yynrhs = yyr2[yyrule]; + int yyi; + unsigned long int yylno = yyrline[yyrule]; + YYFPRINTF (stderr, "Reducing stack by rule %d (line %lu):\n", + yyrule - 1, yylno); + /* The symbols being reduced. */ + for (yyi = 0; yyi < yynrhs; yyi++) + { + YYFPRINTF (stderr, " $%d = ", yyi + 1); + yy_symbol_print (stderr, yyrhs[yyprhs[yyrule] + yyi], + &(yyvsp[(yyi + 1) - (yynrhs)]) + , &(yylsp[(yyi + 1) - (yynrhs)]) , lexer, state); + YYFPRINTF (stderr, "\n"); + } +} + +# define YY_REDUCE_PRINT(Rule) \ +do { \ + if (yydebug) \ + yy_reduce_print (yyvsp, yylsp, Rule, lexer, state); \ +} while (YYID (0)) + +/* Nonzero means print parse trace. It is left uninitialized so that + multiple parsers can coexist. */ +int yydebug; +#else /* !YYDEBUG */ +# define YYDPRINTF(Args) +# define YY_SYMBOL_PRINT(Title, Type, Value, Location) +# define YY_STACK_PRINT(Bottom, Top) +# define YY_REDUCE_PRINT(Rule) +#endif /* !YYDEBUG */ + + +/* YYINITDEPTH -- initial size of the parser's stacks. */ +#ifndef YYINITDEPTH +# define YYINITDEPTH 200 +#endif + +/* YYMAXDEPTH -- maximum size the stacks can grow to (effective only + if the built-in stack extension method is used). + + Do not make this value too large; the results are undefined if + YYSTACK_ALLOC_MAXIMUM < YYSTACK_BYTES (YYMAXDEPTH) + evaluated with infinite-precision integer arithmetic. */ + +#ifndef YYMAXDEPTH +# define YYMAXDEPTH 10000 +#endif + + + +#if YYERROR_VERBOSE + +# ifndef yystrlen +# if defined __GLIBC__ && defined _STRING_H +# define yystrlen strlen +# else +/* Return the length of YYSTR. */ +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static YYSIZE_T +yystrlen (const char *yystr) +#else +static YYSIZE_T +yystrlen (yystr) + const char *yystr; +#endif +{ + YYSIZE_T yylen; + for (yylen = 0; yystr[yylen]; yylen++) + continue; + return yylen; +} +# endif +# endif + +# ifndef yystpcpy +# if defined __GLIBC__ && defined _STRING_H && defined _GNU_SOURCE +# define yystpcpy stpcpy +# else +/* Copy YYSRC to YYDEST, returning the address of the terminating '\0' in + YYDEST. */ +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static char * +yystpcpy (char *yydest, const char *yysrc) +#else +static char * +yystpcpy (yydest, yysrc) + char *yydest; + const char *yysrc; +#endif +{ + char *yyd = yydest; + const char *yys = yysrc; + + while ((*yyd++ = *yys++) != '\0') + continue; + + return yyd - 1; +} +# endif +# endif + +# ifndef yytnamerr +/* Copy to YYRES the contents of YYSTR after stripping away unnecessary + quotes and backslashes, so that it's suitable for yyerror. The + heuristic is that double-quoting is unnecessary unless the string + contains an apostrophe, a comma, or backslash (other than + backslash-backslash). YYSTR is taken from yytname. If YYRES is + null, do not copy; instead, return the length of what the result + would have been. */ +static YYSIZE_T +yytnamerr (char *yyres, const char *yystr) +{ + if (*yystr == '"') + { + YYSIZE_T yyn = 0; + char const *yyp = yystr; + + for (;;) + switch (*++yyp) + { + case '\'': + case ',': + goto do_not_strip_quotes; + + case '\\': + if (*++yyp != '\\') + goto do_not_strip_quotes; + /* Fall through. */ + default: + if (yyres) + yyres[yyn] = *yyp; + yyn++; + break; + + case '"': + if (yyres) + yyres[yyn] = '\0'; + return yyn; + } + do_not_strip_quotes: ; + } + + if (! yyres) + return yystrlen (yystr); + + return yystpcpy (yyres, yystr) - yyres; +} +# endif + +/* Copy into YYRESULT an error message about the unexpected token + YYCHAR while in state YYSTATE. Return the number of bytes copied, + including the terminating null byte. If YYRESULT is null, do not + copy anything; just return the number of bytes that would be + copied. As a special case, return 0 if an ordinary "syntax error" + message will do. Return YYSIZE_MAXIMUM if overflow occurs during + size calculation. */ +static YYSIZE_T +yysyntax_error (char *yyresult, int yystate, int yychar) +{ + int yyn = yypact[yystate]; + + if (! (YYPACT_NINF < yyn && yyn <= YYLAST)) + return 0; + else + { + int yytype = YYTRANSLATE (yychar); + YYSIZE_T yysize0 = yytnamerr (0, yytname[yytype]); + YYSIZE_T yysize = yysize0; + YYSIZE_T yysize1; + int yysize_overflow = 0; + enum { YYERROR_VERBOSE_ARGS_MAXIMUM = 5 }; + char const *yyarg[YYERROR_VERBOSE_ARGS_MAXIMUM]; + int yyx; + +# if 0 + /* This is so xgettext sees the translatable formats that are + constructed on the fly. */ + YY_("syntax error, unexpected %s"); + YY_("syntax error, unexpected %s, expecting %s"); + YY_("syntax error, unexpected %s, expecting %s or %s"); + YY_("syntax error, unexpected %s, expecting %s or %s or %s"); + YY_("syntax error, unexpected %s, expecting %s or %s or %s or %s"); +# endif + char *yyfmt; + char const *yyf; + static char const yyunexpected[] = "syntax error, unexpected %s"; + static char const yyexpecting[] = ", expecting %s"; + static char const yyor[] = " or %s"; + char yyformat[sizeof yyunexpected + + sizeof yyexpecting - 1 + + ((YYERROR_VERBOSE_ARGS_MAXIMUM - 2) + * (sizeof yyor - 1))]; + char const *yyprefix = yyexpecting; + + /* Start YYX at -YYN if negative to avoid negative indexes in + YYCHECK. */ + int yyxbegin = yyn < 0 ? -yyn : 0; + + /* Stay within bounds of both yycheck and yytname. */ + int yychecklim = YYLAST - yyn + 1; + int yyxend = yychecklim < YYNTOKENS ? yychecklim : YYNTOKENS; + int yycount = 1; + + yyarg[0] = yytname[yytype]; + yyfmt = yystpcpy (yyformat, yyunexpected); + + for (yyx = yyxbegin; yyx < yyxend; ++yyx) + if (yycheck[yyx + yyn] == yyx && yyx != YYTERROR) + { + if (yycount == YYERROR_VERBOSE_ARGS_MAXIMUM) + { + yycount = 1; + yysize = yysize0; + yyformat[sizeof yyunexpected - 1] = '\0'; + break; + } + yyarg[yycount++] = yytname[yyx]; + yysize1 = yysize + yytnamerr (0, yytname[yyx]); + yysize_overflow |= (yysize1 < yysize); + yysize = yysize1; + yyfmt = yystpcpy (yyfmt, yyprefix); + yyprefix = yyor; + } + + yyf = YY_(yyformat); + yysize1 = yysize + yystrlen (yyf); + yysize_overflow |= (yysize1 < yysize); + yysize = yysize1; + + if (yysize_overflow) + return YYSIZE_MAXIMUM; + + if (yyresult) + { + /* Avoid sprintf, as that infringes on the user's name space. + Don't have undefined behavior even if the translation + produced a string with the wrong number of "%s"s. */ + char *yyp = yyresult; + int yyi = 0; + while ((*yyp = *yyf) != '\0') + { + if (*yyp == '%' && yyf[1] == 's' && yyi < yycount) + { + yyp += yytnamerr (yyp, yyarg[yyi++]); + yyf += 2; + } + else + { + yyp++; + yyf++; + } + } + } + return yysize; + } +} +#endif /* YYERROR_VERBOSE */ + + +/*-----------------------------------------------. +| Release the memory associated to this symbol. | +`-----------------------------------------------*/ + +/*ARGSUSED*/ +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +static void +yydestruct (const char *yymsg, int yytype, YYSTYPE *yyvaluep, YYLTYPE *yylocationp, parser::PTXLexer& lexer, parser::PTXParser& state) +#else +static void +yydestruct (yymsg, yytype, yyvaluep, yylocationp, lexer, state) + const char *yymsg; + int yytype; + YYSTYPE *yyvaluep; + YYLTYPE *yylocationp; + parser::PTXLexer& lexer; + parser::PTXParser& state; +#endif +{ + YYUSE (yyvaluep); + YYUSE (yylocationp); + YYUSE (lexer); + YYUSE (state); + + if (!yymsg) + yymsg = "Deleting"; + YY_SYMBOL_PRINT (yymsg, yytype, yyvaluep, yylocationp); + + switch (yytype) + { + + default: + break; + } +} + +/* Prevent warnings from -Wmissing-prototypes. */ +#ifdef YYPARSE_PARAM +#if defined __STDC__ || defined __cplusplus +int yyparse (void *YYPARSE_PARAM); +#else +int yyparse (); +#endif +#else /* ! YYPARSE_PARAM */ +#if defined __STDC__ || defined __cplusplus +int yyparse (parser::PTXLexer& lexer, parser::PTXParser& state); +#else +int yyparse (); +#endif +#endif /* ! YYPARSE_PARAM */ + + + + + +/*-------------------------. +| yyparse or yypush_parse. | +`-------------------------*/ + +#ifdef YYPARSE_PARAM +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +int +yyparse (void *YYPARSE_PARAM) +#else +int +yyparse (YYPARSE_PARAM) + void *YYPARSE_PARAM; +#endif +#else /* ! YYPARSE_PARAM */ +#if (defined __STDC__ || defined __C99__FUNC__ \ + || defined __cplusplus || defined _MSC_VER) +int +yyparse (parser::PTXLexer& lexer, parser::PTXParser& state) +#else +int +yyparse (lexer, state) + parser::PTXLexer& lexer; + parser::PTXParser& state; +#endif +#endif +{ +/* The lookahead symbol. */ +int yychar; + +/* The semantic value of the lookahead symbol. */ +YYSTYPE yylval; + +/* Location data for the lookahead symbol. */ +YYLTYPE yylloc; + + /* Number of syntax errors so far. */ + int yynerrs; + + int yystate; + /* Number of tokens to shift before error messages enabled. */ + int yyerrstatus; + + /* The stacks and their tools: + `yyss': related to states. + `yyvs': related to semantic values. + `yyls': related to locations. + + Refer to the stacks thru separate pointers, to allow yyoverflow + to reallocate them elsewhere. */ + + /* The state stack. */ + yytype_int16 yyssa[YYINITDEPTH]; + yytype_int16 *yyss; + yytype_int16 *yyssp; + + /* The semantic value stack. */ + YYSTYPE yyvsa[YYINITDEPTH]; + YYSTYPE *yyvs; + YYSTYPE *yyvsp; + + /* The location stack. */ + YYLTYPE yylsa[YYINITDEPTH]; + YYLTYPE *yyls; + YYLTYPE *yylsp; + + /* The locations where the error started and ended. */ + YYLTYPE yyerror_range[2]; + + YYSIZE_T yystacksize; + + int yyn; + int yyresult; + /* Lookahead token as an internal (translated) token number. */ + int yytoken; + /* The variables used to return semantic value and location from the + action routines. */ + YYSTYPE yyval; + YYLTYPE yyloc; + +#if YYERROR_VERBOSE + /* Buffer for error messages, and its allocated size. */ + char yymsgbuf[128]; + char *yymsg = yymsgbuf; + YYSIZE_T yymsg_alloc = sizeof yymsgbuf; +#endif + +#define YYPOPSTACK(N) (yyvsp -= (N), yyssp -= (N), yylsp -= (N)) + + /* The number of symbols on the RHS of the reduced rule. + Keep to zero when no symbol should be popped. */ + int yylen = 0; + + yytoken = 0; + yyss = yyssa; + yyvs = yyvsa; + yyls = yylsa; + yystacksize = YYINITDEPTH; + + YYDPRINTF ((stderr, "Starting parse\n")); + + yystate = 0; + yyerrstatus = 0; + yynerrs = 0; + yychar = YYEMPTY; /* Cause a token to be read. */ + + /* Initialize stack pointers. + Waste one element of value and location stack + so that they stay on the same level as the state stack. + The wasted elements are never initialized. */ + yyssp = yyss; + yyvsp = yyvs; + yylsp = yyls; + +#if YYLTYPE_IS_TRIVIAL + /* Initialize the default location before parsing starts. */ + yylloc.first_line = yylloc.last_line = 1; + yylloc.first_column = yylloc.last_column = 1; +#endif + + goto yysetstate; + +/*------------------------------------------------------------. +| yynewstate -- Push a new state, which is found in yystate. | +`------------------------------------------------------------*/ + yynewstate: + /* In all cases, when you get here, the value and location stacks + have just been pushed. So pushing a state here evens the stacks. */ + yyssp++; + + yysetstate: + *yyssp = yystate; + + if (yyss + yystacksize - 1 <= yyssp) + { + /* Get the current used size of the three stacks, in elements. */ + YYSIZE_T yysize = yyssp - yyss + 1; + +#ifdef yyoverflow + { + /* Give user a chance to reallocate the stack. Use copies of + these so that the &'s don't force the real ones into + memory. */ + YYSTYPE *yyvs1 = yyvs; + yytype_int16 *yyss1 = yyss; + YYLTYPE *yyls1 = yyls; + + /* Each stack pointer address is followed by the size of the + data in use in that stack, in bytes. This used to be a + conditional around just the two extra args, but that might + be undefined if yyoverflow is a macro. */ + yyoverflow (YY_("memory exhausted"), + &yyss1, yysize * sizeof (*yyssp), + &yyvs1, yysize * sizeof (*yyvsp), + &yyls1, yysize * sizeof (*yylsp), + &yystacksize); + + yyls = yyls1; + yyss = yyss1; + yyvs = yyvs1; + } +#else /* no yyoverflow */ +# ifndef YYSTACK_RELOCATE + goto yyexhaustedlab; +# else + /* Extend the stack our own way. */ + if (YYMAXDEPTH <= yystacksize) + goto yyexhaustedlab; + yystacksize *= 2; + if (YYMAXDEPTH < yystacksize) + yystacksize = YYMAXDEPTH; + + { + yytype_int16 *yyss1 = yyss; + union yyalloc *yyptr = + (union yyalloc *) YYSTACK_ALLOC (YYSTACK_BYTES (yystacksize)); + if (! yyptr) + goto yyexhaustedlab; + YYSTACK_RELOCATE (yyss_alloc, yyss); + YYSTACK_RELOCATE (yyvs_alloc, yyvs); + YYSTACK_RELOCATE (yyls_alloc, yyls); +# undef YYSTACK_RELOCATE + if (yyss1 != yyssa) + YYSTACK_FREE (yyss1); + } +# endif +#endif /* no yyoverflow */ + + yyssp = yyss + yysize - 1; + yyvsp = yyvs + yysize - 1; + yylsp = yyls + yysize - 1; + + YYDPRINTF ((stderr, "Stack size increased to %lu\n", + (unsigned long int) yystacksize)); + + if (yyss + yystacksize - 1 <= yyssp) + YYABORT; + } + + YYDPRINTF ((stderr, "Entering state %d\n", yystate)); + + if (yystate == YYFINAL) + YYACCEPT; + + goto yybackup; + +/*-----------. +| yybackup. | +`-----------*/ +yybackup: + + /* Do appropriate processing given the current state. Read a + lookahead token if we need one and don't already have one. */ + + /* First try to decide what to do without reference to lookahead token. */ + yyn = yypact[yystate]; + if (yyn == YYPACT_NINF) + goto yydefault; + + /* Not known => get a lookahead token if don't already have one. */ + + /* YYCHAR is either YYEMPTY or YYEOF or a valid lookahead symbol. */ + if (yychar == YYEMPTY) + { + YYDPRINTF ((stderr, "Reading a token: ")); + yychar = YYLEX; + } + + if (yychar <= YYEOF) + { + yychar = yytoken = YYEOF; + YYDPRINTF ((stderr, "Now at end of input.\n")); + } + else + { + yytoken = YYTRANSLATE (yychar); + YY_SYMBOL_PRINT ("Next token is", yytoken, &yylval, &yylloc); + } + + /* If the proper action on seeing token YYTOKEN is to reduce or to + detect an error, take that action. */ + yyn += yytoken; + if (yyn < 0 || YYLAST < yyn || yycheck[yyn] != yytoken) + goto yydefault; + yyn = yytable[yyn]; + if (yyn <= 0) + { + if (yyn == 0 || yyn == YYTABLE_NINF) + goto yyerrlab; + yyn = -yyn; + goto yyreduce; + } + + /* Count tokens shifted since error; after three, turn off error + status. */ + if (yyerrstatus) + yyerrstatus--; + + /* Shift the lookahead token. */ + YY_SYMBOL_PRINT ("Shifting", yytoken, &yylval, &yylloc); + + /* Discard the shifted token. */ + yychar = YYEMPTY; + + yystate = yyn; + *++yyvsp = yylval; + *++yylsp = yylloc; + goto yynewstate; + + +/*-----------------------------------------------------------. +| yydefault -- do the default action for the current state. | +`-----------------------------------------------------------*/ +yydefault: + yyn = yydefact[yystate]; + if (yyn == 0) + goto yyerrlab; + goto yyreduce; + + +/*-----------------------------. +| yyreduce -- Do a reduction. | +`-----------------------------*/ +yyreduce: + /* yyn is the number of a rule to reduce with. */ + yylen = yyr2[yyn]; + + /* If YYLEN is nonzero, implement the default value of the action: + `$$ = $1'. + + Otherwise, the following line sets YYVAL to garbage. + This behavior is undocumented and Bison + users should not rely upon it. Assigning to YYVAL + unconditionally makes the parser a bit smaller, and it avoids a + GCC warning that YYVAL may be used uninitialized. */ + yyval = yyvsp[1-yylen]; + + /* Default location. */ + YYLLOC_DEFAULT (yyloc, (yylsp - yylen), yylen); + YY_REDUCE_PRINT (yyn); + switch (yyn) + { + case 3: + +/* Line 1455 of yacc.c */ +#line 75 "ptxgrammar.yy" + { + std::cerr << "Done reading PTX \n" << std::endl; + state.printHeader(); +;} + break; + + case 4: + +/* Line 1455 of yacc.c */ +#line 81 "ptxgrammar.yy" + { std::cerr << "Reading PTX version " << (yyvsp[(2) - (2)].fvalue) << std::endl; ;} + break; + + case 5: + +/* Line 1455 of yacc.c */ +#line 83 "ptxgrammar.yy" + { std::cerr << "Target " << (yyvsp[(2) - (2)].svalue) << std::endl; ;} + break; + + case 6: + +/* Line 1455 of yacc.c */ +#line 85 "ptxgrammar.yy" + { std::cerr << "Address_Size " << (yyvsp[(2) - (2)].ivalue) << std::endl; ;} + break; + + case 21: + +/* Line 1455 of yacc.c */ +#line 94 "ptxgrammar.yy" + { state.dataTypeId((yyvsp[(1) - (1)].ivalue)); ;} + break; + + case 44: + +/* Line 1455 of yacc.c */ +#line 117 "ptxgrammar.yy" + { (yyval.ivalue) = (yyvsp[(2) - (3)].ivalue); state.arrayDimensions((yyvsp[(2) - (3)].ivalue)); ;} + break; + + case 47: + +/* Line 1455 of yacc.c */ +#line 123 "ptxgrammar.yy" + { strcpy((yyval.svalue), (yyvsp[(1) - (1)].svalue)); state.identifier((yyvsp[(1) - (1)].svalue)); ;} + break; + + case 49: + +/* Line 1455 of yacc.c */ +#line 126 "ptxgrammar.yy" + {(yyval.ivalue) = (yyvsp[(2) - (2)].ivalue); state.alignment((yyvsp[(2) - (2)].ivalue));;} + break; + + case 50: + +/* Line 1455 of yacc.c */ +#line 127 "ptxgrammar.yy" + { state.alignment(0); ;} + break; + + case 52: + +/* Line 1455 of yacc.c */ +#line 131 "ptxgrammar.yy" + { + state.argumentDeclaration((yylsp[(1) - (4)])); +;} + break; + + case 53: + +/* Line 1455 of yacc.c */ +#line 136 "ptxgrammar.yy" + { state.argumentListBegin((yylsp[(1) - (1)])); ;} + break; + + case 54: + +/* Line 1455 of yacc.c */ +#line 137 "ptxgrammar.yy" + {state.argumentListEnd((yylsp[(1) - (1)])); ;} + break; + + case 59: + +/* Line 1455 of yacc.c */ +#line 144 "ptxgrammar.yy" + { + state.visibleEntryDeclaration((yyvsp[(3) - (4)].svalue), (yylsp[(1) - (4)])); +;} + break; + + case 60: + +/* Line 1455 of yacc.c */ +#line 148 "ptxgrammar.yy" + { state.returnArgumentListBegin((yylsp[(1) - (1)])); ;} + break; + + case 61: + +/* Line 1455 of yacc.c */ +#line 149 "ptxgrammar.yy" + {state.returnArgumentListEnd((yylsp[(1) - (1)])); ;} + break; + + case 65: + +/* Line 1455 of yacc.c */ +#line 153 "ptxgrammar.yy" + { + state.visibleFunctionDeclaration((yyvsp[(4) - (5)].svalue), (yylsp[(1) - (5)])); +;} + break; + + case 66: + +/* Line 1455 of yacc.c */ +#line 159 "ptxgrammar.yy" + { + state.visibleInitializableDeclaration((yyvsp[(4) - (5)].svalue),(yylsp[(1) - (5)])); +;} + break; + + + +/* Line 1455 of yacc.c */ +#line 1664 "ptxgrammar.cc" + default: break; + } + YY_SYMBOL_PRINT ("-> $$ =", yyr1[yyn], &yyval, &yyloc); + + YYPOPSTACK (yylen); + yylen = 0; + YY_STACK_PRINT (yyss, yyssp); + + *++yyvsp = yyval; + *++yylsp = yyloc; + + /* Now `shift' the result of the reduction. Determine what state + that goes to, based on the state we popped back to and the rule + number reduced by. */ + + yyn = yyr1[yyn]; + + yystate = yypgoto[yyn - YYNTOKENS] + *yyssp; + if (0 <= yystate && yystate <= YYLAST && yycheck[yystate] == *yyssp) + yystate = yytable[yystate]; + else + yystate = yydefgoto[yyn - YYNTOKENS]; + + goto yynewstate; + + +/*------------------------------------. +| yyerrlab -- here on detecting error | +`------------------------------------*/ +yyerrlab: + /* If not already recovering from an error, report this error. */ + if (!yyerrstatus) + { + ++yynerrs; +#if ! YYERROR_VERBOSE + yyerror (&yylloc, lexer, state, YY_("syntax error")); +#else + { + YYSIZE_T yysize = yysyntax_error (0, yystate, yychar); + if (yymsg_alloc < yysize && yymsg_alloc < YYSTACK_ALLOC_MAXIMUM) + { + YYSIZE_T yyalloc = 2 * yysize; + if (! (yysize <= yyalloc && yyalloc <= YYSTACK_ALLOC_MAXIMUM)) + yyalloc = YYSTACK_ALLOC_MAXIMUM; + if (yymsg != yymsgbuf) + YYSTACK_FREE (yymsg); + yymsg = (char *) YYSTACK_ALLOC (yyalloc); + if (yymsg) + yymsg_alloc = yyalloc; + else + { + yymsg = yymsgbuf; + yymsg_alloc = sizeof yymsgbuf; + } + } + + if (0 < yysize && yysize <= yymsg_alloc) + { + (void) yysyntax_error (yymsg, yystate, yychar); + yyerror (&yylloc, lexer, state, yymsg); + } + else + { + yyerror (&yylloc, lexer, state, YY_("syntax error")); + if (yysize != 0) + goto yyexhaustedlab; + } + } +#endif + } + + yyerror_range[0] = yylloc; + + if (yyerrstatus == 3) + { + /* If just tried and failed to reuse lookahead token after an + error, discard it. */ + + if (yychar <= YYEOF) + { + /* Return failure if at end of input. */ + if (yychar == YYEOF) + YYABORT; + } + else + { + yydestruct ("Error: discarding", + yytoken, &yylval, &yylloc, lexer, state); + yychar = YYEMPTY; + } + } + + /* Else will try to reuse lookahead token after shifting the error + token. */ + goto yyerrlab1; + + +/*---------------------------------------------------. +| yyerrorlab -- error raised explicitly by YYERROR. | +`---------------------------------------------------*/ +yyerrorlab: + + /* Pacify compilers like GCC when the user code never invokes + YYERROR and the label yyerrorlab therefore never appears in user + code. */ + if (/*CONSTCOND*/ 0) + goto yyerrorlab; + + yyerror_range[0] = yylsp[1-yylen]; + /* Do not reclaim the symbols of the rule which action triggered + this YYERROR. */ + YYPOPSTACK (yylen); + yylen = 0; + YY_STACK_PRINT (yyss, yyssp); + yystate = *yyssp; + goto yyerrlab1; + + +/*-------------------------------------------------------------. +| yyerrlab1 -- common code for both syntax error and YYERROR. | +`-------------------------------------------------------------*/ +yyerrlab1: + yyerrstatus = 3; /* Each real token shifted decrements this. */ + + for (;;) + { + yyn = yypact[yystate]; + if (yyn != YYPACT_NINF) + { + yyn += YYTERROR; + if (0 <= yyn && yyn <= YYLAST && yycheck[yyn] == YYTERROR) + { + yyn = yytable[yyn]; + if (0 < yyn) + break; + } + } + + /* Pop the current state because it cannot handle the error token. */ + if (yyssp == yyss) + YYABORT; + + yyerror_range[0] = *yylsp; + yydestruct ("Error: popping", + yystos[yystate], yyvsp, yylsp, lexer, state); + YYPOPSTACK (1); + yystate = *yyssp; + YY_STACK_PRINT (yyss, yyssp); + } + + *++yyvsp = yylval; + + yyerror_range[1] = yylloc; + /* Using YYLLOC is tempting, but would change the location of + the lookahead. YYLOC is available though. */ + YYLLOC_DEFAULT (yyloc, (yyerror_range - 1), 2); + *++yylsp = yyloc; + + /* Shift the error token. */ + YY_SYMBOL_PRINT ("Shifting", yystos[yyn], yyvsp, yylsp); + + yystate = yyn; + goto yynewstate; + + +/*-------------------------------------. +| yyacceptlab -- YYACCEPT comes here. | +`-------------------------------------*/ +yyacceptlab: + yyresult = 0; + goto yyreturn; + +/*-----------------------------------. +| yyabortlab -- YYABORT comes here. | +`-----------------------------------*/ +yyabortlab: + yyresult = 1; + goto yyreturn; + +#if !defined(yyoverflow) || YYERROR_VERBOSE +/*-------------------------------------------------. +| yyexhaustedlab -- memory exhaustion comes here. | +`-------------------------------------------------*/ +yyexhaustedlab: + yyerror (&yylloc, lexer, state, YY_("memory exhausted")); + yyresult = 2; + /* Fall through. */ +#endif + +yyreturn: + if (yychar != YYEMPTY) + yydestruct ("Cleanup: discarding lookahead", + yytoken, &yylval, &yylloc, lexer, state); + /* Do not reclaim the symbols of the rule which action triggered + this YYABORT or YYACCEPT. */ + YYPOPSTACK (yylen); + YY_STACK_PRINT (yyss, yyssp); + while (yyssp != yyss) + { + yydestruct ("Cleanup: popping", + yystos[*yyssp], yyvsp, yylsp, lexer, state); + YYPOPSTACK (1); + } +#ifndef yyoverflow + if (yyss != yyssa) + YYSTACK_FREE (yyss); +#endif +#if YYERROR_VERBOSE + if (yymsg != yymsgbuf) + YYSTACK_FREE (yymsg); +#endif + /* Make sure YYID is used. */ + return YYID (yyresult); +} + + + +/* Line 1675 of yacc.c */ +#line 164 "ptxgrammar.yy" + + +int yylex( YYSTYPE* token, YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state ) +{ + lexer.yylval = token; + + int tokenValue = lexer.yylexPosition(); + location->first_line = lexer.lineno(); + location->first_column = lexer.column; + +#if 0 + report( " Lexer (" << location->first_line << "," + << location->first_column + << "): " << parser::PTXLexer::toString( tokenValue ) << " \"" + << lexer.YYText() << "\""); +#endif + + return tokenValue; +} + +static std::string toString( YYLTYPE& location, parser::PTXParser& state ) +{ + std::stringstream stream; + stream +#if 0 + << state.fileName +#else + << "ptx " +#endif + << " (" << location.first_line << ", " + << location.first_column << "): "; + return stream.str(); +} + +void yyerror( YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state, char const* message ) +{ + std::stringstream stream; + stream << toString( *location, state ) + << " " << message; + fprintf(stderr, "--ERROR-- %s %s \n", toString(*location, state).c_str(), message); + assert(0); +} + +} + diff --git a/examples_ptx/ptxcc/ptxgrammar.yy b/examples_ptx/ptxcc/ptxgrammar.yy new file mode 100644 index 00000000..009ba140 --- /dev/null +++ b/examples_ptx/ptxcc/ptxgrammar.yy @@ -0,0 +1,210 @@ +%locations + +%{ + #include + #include "PTXParser.h" + #include "PTXLexer.h" + #include + #include + #include + #include + + #define YYERROR_VERBOSE 1 + + #ifdef REPORT_BASE + #undef REPORT_BASE + #endif + + #define REPORT_BASE 0 + + namespace ptx + { + + int yylex( YYSTYPE* token, YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state ); + void yyerror( YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state, char const* message ); + + std::string yyTypeToString( int ); + +%} + +%union +{ + char svalue[1024]; + double fvalue; + int ivalue; + unsigned int uvalue; +} + +%parse-param {parser::PTXLexer& lexer} +%parse-param {parser::PTXParser& state} +%lex-param {parser::PTXLexer& lexer} +%lex-param {parser::PTXParser& state} +%pure-parser + +// define the constant-string tokens: +%token TOKEN_VERSION TOKEN_TARGET TOKEN_ADDRESS_SIZE +%token TOKEN_VISIBLE TOKEN_FUNC TOKEN_ENTRY +%token TOKEN_PARAM TOKEN_ALIGN +%token TOKEN_GLOBAL +%token TOKEN_B8 TOKEN_B16 TOKEN_B32 TOKEN_B64 +%token TOKEN_U8 TOKEN_U16 TOKEN_U32 TOKEN_U64 +%token TOKEN_S8 TOKEN_S16 TOKEN_S32 TOKEN_S64 +%token TOKEN_F32 TOKEN_F64 + +// define the "terminal symbol" token types I'm going to use (in CAPS +// by convention), and associate each with a field of the union: +%token TOKEN_INT +%token TOKEN_FLOAT +%token TOKEN_STRING + +%type identifier +%type arrayDimensionSet +%type alignment + +%start ptxsource + +%% +// the first rule defined is the highest-level rule, which in our +// case is just the concept of a whole "snazzle file": +ptxsource: + header ptxbody; + +header: + version target address_size +{ + std::cerr << "Done reading PTX \n" << std::endl; + state.printHeader(); +}; + +version: + TOKEN_VERSION TOKEN_FLOAT { std::cerr << "Reading PTX version " << $2 << std::endl; }; +target: + TOKEN_TARGET TOKEN_STRING { std::cerr << "Target " << $2 << std::endl; }; +address_size: + TOKEN_ADDRESS_SIZE TOKEN_INT { std::cerr << "Address_Size " << $2 << std::endl; }; + + +dataTypeId : + TOKEN_U8 | TOKEN_U16 | TOKEN_U32 | TOKEN_U64 + | TOKEN_S8 | TOKEN_S16 | TOKEN_S32 | TOKEN_S64 + | TOKEN_B8 | TOKEN_B16 | TOKEN_B32 | TOKEN_B64 + | TOKEN_F32 | TOKEN_F64; + +dataType: dataTypeId { state.dataTypeId($1); } + +anytoken: + TOKEN_ALIGN +| TOKEN_PARAM +| dataTypeId +| TOKEN_STRING | TOKEN_FLOAT | TOKEN_INT +| TOKEN_FUNC | TOKEN_ENTRY +| TOKEN_GLOBAL +| '[' +| ']' +| '(' +| ')' +| ','; + +ptxbody: + ptxbody visibleFunctionDeclaration | visibleFunctionDeclaration + | ptxbody visibleEntryDeclaration| visibleEntryDeclaration + | ptxbody visibleInitializableDeclaration| visibleInitializableDeclaration + | ptxbody anytoken | anytoken; + + + +arrayDimensionSet : '[' TOKEN_INT ']' { $$ = $2; state.arrayDimensions($2); } +// arrayDimensionSet : arrayDimensionSet '[' TOKEN_INT ']' { $$ = $2; } +// arrayDimensionSet : '[' ']' { $$ = 0; } +arrayDimensions : /* empty string */; +arrayDimensions : arrayDimensionSet; + +identifier: TOKEN_STRING { strcpy($$, $1); state.identifier($1); } +parameter : TOKEN_PARAM; + +alignment : TOKEN_ALIGN TOKEN_INT {$$ = $2; state.alignment($2);} +addressableVariablePrefix : dataType { state.alignment(0); } +addressableVariablePrefix : alignment dataType; + +argumentDeclaration : parameter addressableVariablePrefix identifier arrayDimensions +{ + state.argumentDeclaration(@1); +} + + +argumentListBegin : '(' { state.argumentListBegin(@1); }; +argumentListEnd : ')' {state.argumentListEnd(@1); }; +argumentListBody : argumentDeclaration; +argumentListBody : /* empty string */; +argumentListBody : argumentListBody ',' argumentDeclaration; +argumentList: argumentListBegin argumentListBody argumentListEnd; + +visibleEntryDeclaration: TOKEN_VISIBLE TOKEN_ENTRY identifier argumentList +{ + state.visibleEntryDeclaration($3, @1); +}; + +returnArgumentListBegin : '(' { state.returnArgumentListBegin(@1); } +returnArgumentListEnd : ')' {state.returnArgumentListEnd(@1); } +returnArgumentList : returnArgumentListBegin argumentListBody returnArgumentListEnd; +optionalReturnArgumentList : returnArgumentList | /* empty string */; +visibleFunctionDeclaration: TOKEN_VISIBLE TOKEN_FUNC optionalReturnArgumentList identifier argumentList +{ + state.visibleFunctionDeclaration($4, @1); +}; + +visibleInitializableDeclaration : + TOKEN_VISIBLE TOKEN_GLOBAL addressableVariablePrefix identifier arrayDimensionSet +{ + state.visibleInitializableDeclaration($4,@1); +} + + +%% + +int yylex( YYSTYPE* token, YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state ) +{ + lexer.yylval = token; + + int tokenValue = lexer.yylexPosition(); + location->first_line = lexer.lineno(); + location->first_column = lexer.column; + +#if 0 + report( " Lexer (" << location->first_line << "," + << location->first_column + << "): " << parser::PTXLexer::toString( tokenValue ) << " \"" + << lexer.YYText() << "\""); +#endif + + return tokenValue; +} + +static std::string toString( YYLTYPE& location, parser::PTXParser& state ) +{ + std::stringstream stream; + stream +#if 0 + << state.fileName +#else + << "ptx " +#endif + << " (" << location.first_line << ", " + << location.first_column << "): "; + return stream.str(); +} + +void yyerror( YYLTYPE* location, parser::PTXLexer& lexer, + parser::PTXParser& state, char const* message ) +{ + std::stringstream stream; + stream << toString( *location, state ) + << " " << message; + fprintf(stderr, "--ERROR-- %s %s \n", toString(*location, state).c_str(), message); + assert(0); +} + +} diff --git a/examples_ptx/ptxgen/Makefile b/examples_ptx/ptxgen/Makefile new file mode 100644 index 00000000..f1aad5b6 --- /dev/null +++ b/examples_ptx/ptxgen/Makefile @@ -0,0 +1,15 @@ +CUDA55 = /usr/local/cuda-5.5 +CUDA60 = /usr/local/cuda-6.0 + +all: ptxgen55 ptxgen60 + +ptxgen60: ptxgen.c + gcc -O3 -o ptxgen60 ptxgen.c -L$(CUDA60)/nvvm/lib64 -lnvvm -I$(CUDA60)/nvvm/include -I$(CUDA60)/include -DLIBDEVICE_MAJOR_VERSION=1 -DLIBDEVICE_MINOR_VERSION=0 -DLIBNVVM_HOME=$(CUDA60)/nvvm -Wl,-rpath,$(CUDA60)/nvvm/lib64 + + +ptxgen55: ptxgen.c + gcc -O3 -o ptxgen55 ptxgen.c -L$(CUDA55)/nvvm/lib64 -lnvvm -I$(CUDA55)/nvvm/include -I$(CUDA55)/include -DLIBDEVICE_MAJOR_VERSION=1 -DLIBDEVICE_MINOR_VERSION=0 -DLIBNVVM_HOME=$(CUDA55)/nvvm -Wl,-rpath,$(CUDA55)/nvvm/lib64 + + +clean: + /bin/rm -f ptxgen55 ptxgen60 diff --git a/examples_ptx/ptxgen/ptxgen.c b/examples_ptx/ptxgen/ptxgen.c new file mode 100644 index 00000000..c91c748e --- /dev/null +++ b/examples_ptx/ptxgen/ptxgen.c @@ -0,0 +1,298 @@ +/* + * Copyright 1993-2012 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +#include + +#include +#include +#include + +/* Two levels of indirection to stringify LIBDEVICE_MAJOR_VERSION and + * LIBDEVICE_MINOR_VERSION correctly. */ +#define getLibDeviceNameForArch(ARCH) \ + _getLibDeviceNameForArch(ARCH, \ + LIBDEVICE_MAJOR_VERSION, \ + LIBDEVICE_MINOR_VERSION) +#define _getLibDeviceNameForArch(ARCH, MAJOR, MINOR) \ + __getLibDeviceNameForArch(ARCH, MAJOR, MINOR) +#define __getLibDeviceNameForArch(ARCH, MAJOR, MINOR) \ + ("/libdevice/libdevice.compute_" #ARCH "." #MAJOR #MINOR ".bc") + +#define getLibnvvmHome _getLibnvvmHome(LIBNVVM_HOME) +#define _getLibnvvmHome(NVVM_HOME) __getLibnvvmHome(NVVM_HOME) +#define __getLibnvvmHome(NVVM_HOME) (#NVVM_HOME) + +typedef struct stat Stat; + +typedef enum { + PTXGEN_SUCCESS = 0x0000, + PTXGEN_FILE_IO_ERROR = 0x0001, + PTXGEN_BAD_ALLOC_ERROR = 0x0002, + PTXGEN_LIBNVVM_COMPILATION_ERROR = 0x0004, + PTXGEN_LIBNVVM_ERROR = 0x0008, + PTXGEN_INVALID_USAGE = 0x0010, + PTXGEN_LIBNVVM_HOME_UNDEFINED = 0x0020, + PTXGEN_LIBNVVM_VERIFICATION_ERROR = 0x0040 +} PTXGENStatus; + +static PTXGENStatus getLibDeviceName(int computeArch, char **buffer) +{ + const char *libnvvmPath = getLibnvvmHome; + const char *libdevice = NULL; + + if (libnvvmPath == NULL) { + fprintf(stderr, "The environment variable LIBNVVM_HOME undefined\n"); + return PTXGEN_LIBNVVM_HOME_UNDEFINED; + } + + /* Use libdevice for compute_20, if the target is not compute_20, compute_30, + * or compute_35. */ + switch (computeArch) { + default: + libdevice = getLibDeviceNameForArch(20); + break; + case 30: + libdevice = getLibDeviceNameForArch(30); + break; + case 35: + libdevice = getLibDeviceNameForArch(35); + break; + } + + *buffer = (char *) malloc(strlen(libnvvmPath) + strlen(libdevice) + 1); + if (*buffer == NULL) { + fprintf(stderr, "Failed to allocate memory\n"); + return PTXGEN_BAD_ALLOC_ERROR; + } + + /* Concatenate libnvvmPath and name. */ + *buffer = strcat(strcpy(*buffer, libnvvmPath), libdevice); + + return PTXGEN_SUCCESS; +} + +static PTXGENStatus addFileToProgram(const char *filename, nvvmProgram prog) +{ + char *buffer; + size_t size; + Stat fileStat; + + /* Open the input file. */ + FILE *f = fopen(filename, "rb"); + if (f == NULL) { + fprintf(stderr, "Failed to open %s\n", filename); + return PTXGEN_FILE_IO_ERROR; + } + + /* Allocate buffer for the input. */ + fstat(fileno(f), &fileStat); + buffer = (char *) malloc(fileStat.st_size); + if (buffer == NULL) { + fprintf(stderr, "Failed to allocate memory\n"); + return PTXGEN_BAD_ALLOC_ERROR; + } + size = fread(buffer, 1, fileStat.st_size, f); + if (ferror(f)) { + fprintf(stderr, "Failed to read %s\n", filename); + fclose(f); + free(buffer); + return PTXGEN_FILE_IO_ERROR; + } + fclose(f); + + if (nvvmAddModuleToProgram(prog, buffer, size, filename) != NVVM_SUCCESS) { + fprintf(stderr, + "Failed to add the module %s to the compilation unit\n", + filename); + free(buffer); + return PTXGEN_LIBNVVM_ERROR; + } + + free(buffer); + return PTXGEN_SUCCESS; +} + +static PTXGENStatus generatePTX(int numOptions, const char **options, + int numFilenames, const char **filenames, + int computeArch) +{ + PTXGENStatus status; + nvvmProgram prog; + char *libDeviceName; + int i; + + /* Create the compiliation unit. */ + if (nvvmCreateProgram(&prog) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to create the compilation unit\n"); + return PTXGEN_LIBNVVM_ERROR; + } + + /* Add libdevice. */ + status = getLibDeviceName(computeArch, &libDeviceName); + if (status != PTXGEN_SUCCESS) { + nvvmDestroyProgram(&prog); + return status; + } + status = addFileToProgram(libDeviceName, prog); + free(libDeviceName); + if (status != PTXGEN_SUCCESS) { + nvvmDestroyProgram(&prog); + return status; + } + + /* Add the module to the compilation unit. */ + for (i = 0; i < numFilenames; ++i) { + status = addFileToProgram(filenames[i], prog); + if (status != PTXGEN_SUCCESS) { + nvvmDestroyProgram(&prog); + return status; + } + } + + /* Verify the compilation unit. */ + if (nvvmVerifyProgram(prog, numOptions, options) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to verify the compilation unit\n"); + status |= PTXGEN_LIBNVVM_VERIFICATION_ERROR; + } + + /* Print warnings and errors. */ + { + size_t logSize; + char *log; + if (nvvmGetProgramLogSize(prog, &logSize) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to get the compilation log size\n"); + status |= PTXGEN_LIBNVVM_ERROR; + } else { + log = (char *) malloc(logSize); + if (log == NULL) { + fprintf(stderr, "Failed to allocate memory\n"); + status |= PTXGEN_BAD_ALLOC_ERROR; + } else if (nvvmGetProgramLog(prog, log) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to get the compilation log\n"); + status |= PTXGEN_LIBNVVM_ERROR; + } else { + fprintf(stderr, "%s\n", log); + } + free(log); + } + } + + if (status & PTXGEN_LIBNVVM_VERIFICATION_ERROR) { + nvvmDestroyProgram(&prog); + return status; + } + + /* Compile the compilation unit. */ + if (nvvmCompileProgram(prog, numOptions, options) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to generate PTX from the compilation unit\n"); + status |= PTXGEN_LIBNVVM_COMPILATION_ERROR; + } else { + size_t ptxSize; + char *ptx; + if (nvvmGetCompiledResultSize(prog, &ptxSize) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to get the PTX output size\n"); + status |= PTXGEN_LIBNVVM_ERROR; + } else { + ptx = (char *) malloc(ptxSize); + if (ptx == NULL) { + fprintf(stderr, "Failed to allocate memory\n"); + status |= PTXGEN_BAD_ALLOC_ERROR; + } else if (nvvmGetCompiledResult(prog, ptx) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to get the PTX output\n"); + status |= PTXGEN_LIBNVVM_ERROR; + } else { + fprintf(stdout, "%s\n", ptx); + } + free(ptx); + } + } + + /* Print warnings and errors. */ + { + size_t logSize; + char *log; + if (nvvmGetProgramLogSize(prog, &logSize) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to get the compilation log size\n"); + status |= PTXGEN_LIBNVVM_ERROR; + } else { + log = (char *) malloc(logSize); + if (log == NULL) { + fprintf(stderr, "Failed to allocate memory\n"); + status |= PTXGEN_BAD_ALLOC_ERROR; + } else if (nvvmGetProgramLog(prog, log) != NVVM_SUCCESS) { + fprintf(stderr, "Failed to get the compilation log\n"); + status |= PTXGEN_LIBNVVM_ERROR; + } else { + fprintf(stderr, "%s\n", log); + } + free(log); + } + } + + /* Release the resources. */ + nvvmDestroyProgram(&prog); + + return PTXGEN_SUCCESS; +} + +static void showUsage() +{ + fprintf(stderr,"Usage: ptxgen [OPTION]... [FILE]...\n" + " [FILE] could be a .bc file or a .ll file\n"); +} + +int main(int argc, char *argv[]) +{ + PTXGENStatus status = PTXGEN_SUCCESS; + int numOptions = 0; + char **options = NULL; + int numFilenames = 0; + char **filenames = NULL; + int computeArch = 35; + int i; + + /* Process the command-line arguments to extract the libnvvm options and the + * input file names. */ + if (argc == 1) { + showUsage(); + return PTXGEN_INVALID_USAGE; + } + + options = (char **) malloc((argc ) * sizeof (char *)); + filenames = (char **) malloc((argc - 1) * sizeof (char *)); + + for (i = 1; i < argc; ++i) { + if (argv[i][0] == '-') { + options[numOptions] = argv[i]; + ++numOptions; + } else { + filenames[numFilenames] = argv[i]; + ++numFilenames; + } + } + options[numOptions] = "-arch=compute_35"; + numOptions++; + + if (numFilenames == 0) { + /* If no input filename is found, then show the usage. */ + showUsage(); + status = PTXGEN_INVALID_USAGE; + } else { + /* Run libnvvm to generate PTX. */ + status = generatePTX(numOptions, (const char **) options, + numFilenames, (const char **) filenames, + computeArch); + } + + free(options); + free(filenames); + return status; +}