added .ptx->.o genreation routines. experimental
This commit is contained in:
Evghenii
26
examples_ptx/ptxcc/Makefile
Normal file
26
examples_ptx/ptxcc/Makefile
Normal file
@@ -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
|
||||
|
||||
40
examples_ptx/ptxcc/PTXLexer.h
Normal file
40
examples_ptx/ptxcc/PTXLexer.h
Normal file
@@ -0,0 +1,40 @@
|
||||
#pragma once
|
||||
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
|
||||
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;
|
||||
}
|
||||
|
||||
};
|
||||
}
|
||||
236
examples_ptx/ptxcc/PTXParser.h
Normal file
236
examples_ptx/ptxcc/PTXParser.h
Normal file
@@ -0,0 +1,236 @@
|
||||
#pragma once
|
||||
|
||||
#undef yyFlexLexer
|
||||
#define yyFlexLexer ptxFlexLexer
|
||||
#include <FlexLexer.h>
|
||||
|
||||
|
||||
#include "PTXLexer.h"
|
||||
|
||||
#include <vector>
|
||||
#include <sstream>
|
||||
#include <string>
|
||||
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<argument_t> argumentList, returnArgumentList;
|
||||
std::vector<int> arrayDimensionsList;
|
||||
|
||||
public:
|
||||
PTXParser(std::ostream &_out) : out(_out)
|
||||
{
|
||||
isArgumentList = isReturnArgumentList = false;
|
||||
_alignment = 1;
|
||||
}
|
||||
|
||||
void printHeader()
|
||||
{
|
||||
std::stringstream s;
|
||||
s << "template<int N> struct __align__(N) b8_t { unsigned char _v[N]; __device__ b8_t() {}; __device__ b8_t (const int value) {}}; \n";
|
||||
s << "template<int N> 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<"<<dim<<"> ";
|
||||
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<"<<dim<<"> ";
|
||||
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();
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
24
examples_ptx/ptxcc/dePTX.cpp
Normal file
24
examples_ptx/ptxcc/dePTX.cpp
Normal file
@@ -0,0 +1,24 @@
|
||||
#include <cstdio>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <cassert>
|
||||
#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());
|
||||
|
||||
}
|
||||
76
examples_ptx/ptxcc/ptx.ll
Normal file
76
examples_ptx/ptxcc/ptx.ll
Normal file
@@ -0,0 +1,76 @@
|
||||
%option yylineno
|
||||
%option noyywrap
|
||||
%option yyclass="parser::PTXLexer"
|
||||
%option prefix="ptx"
|
||||
%option c++
|
||||
|
||||
%{
|
||||
#include "PTXLexer.h"
|
||||
#include <cassert>
|
||||
#include <sstream>
|
||||
#include <cstring>
|
||||
#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
|
||||
18
examples_ptx/ptxcc/ptxcc
Executable file
18
examples_ptx/ptxcc/ptxcc
Executable file
@@ -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
|
||||
1929
examples_ptx/ptxcc/ptxgrammar.cpp
Normal file
1929
examples_ptx/ptxcc/ptxgrammar.cpp
Normal file
File diff suppressed because it is too large
Load Diff
210
examples_ptx/ptxcc/ptxgrammar.yy
Normal file
210
examples_ptx/ptxcc/ptxgrammar.yy
Normal file
@@ -0,0 +1,210 @@
|
||||
%locations
|
||||
|
||||
%{
|
||||
#include <iostream>
|
||||
#include "PTXParser.h"
|
||||
#include "PTXLexer.h"
|
||||
#include <cassert>
|
||||
#include <cstring>
|
||||
#include <sstream>
|
||||
#include <cstdio>
|
||||
|
||||
#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<ivalue> TOKEN_B8 TOKEN_B16 TOKEN_B32 TOKEN_B64
|
||||
%token<ivalue> TOKEN_U8 TOKEN_U16 TOKEN_U32 TOKEN_U64
|
||||
%token<ivalue> TOKEN_S8 TOKEN_S16 TOKEN_S32 TOKEN_S64
|
||||
%token<ivalue> 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 <ivalue> TOKEN_INT
|
||||
%token <fvalue> TOKEN_FLOAT
|
||||
%token <svalue> TOKEN_STRING
|
||||
|
||||
%type<svalue> identifier
|
||||
%type<ivalue> arrayDimensionSet
|
||||
%type<ivalue> 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($<ivalue>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($<ivalue>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($<ivalue>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($<svalue>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($<svalue>4, @1);
|
||||
};
|
||||
|
||||
visibleInitializableDeclaration :
|
||||
TOKEN_VISIBLE TOKEN_GLOBAL addressableVariablePrefix identifier arrayDimensionSet
|
||||
{
|
||||
state.visibleInitializableDeclaration($<svalue>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);
|
||||
}
|
||||
|
||||
}
|
||||
Reference in New Issue
Block a user