diff --git a/ocelot/include/ocelot/parser/PTXParser.h b/ocelot/include/ocelot/parser/PTXParser.h index 1527b55c..a017de5d 100644 --- a/ocelot/include/ocelot/parser/PTXParser.h +++ b/ocelot/include/ocelot/parser/PTXParser.h @@ -150,7 +150,7 @@ namespace parser void maxnreg( unsigned int regs ); void maxntid( unsigned int tidx, unsigned int tidy = 1024, unsigned int tidz = 1024 ); - void ctapersm( int target, unsigned int ctas ); + void ctapersm( const char* target, unsigned int ctas ); void maxnctapersm( unsigned int ctas ); void maxnctapersm(); void minnctapersm( unsigned int ctas ); @@ -171,6 +171,7 @@ namespace parser void singleList( float value ); void singleList1( float value ); void targetElement( int token ); + void targetElement( const char* target ); void target(); void noAddressSpace(); void addressSpace( int token ); diff --git a/ocelot/src/ir/PTXInstruction.cpp b/ocelot/src/ir/PTXInstruction.cpp index 099a0493..c42c7506 100644 --- a/ocelot/src/ir/PTXInstruction.cpp +++ b/ocelot/src/ir/PTXInstruction.cpp @@ -667,12 +667,14 @@ std::string ir::PTXInstruction::valid() const { return "operand 2 type " + PTXOperand::toString( a.type ) + " cannot be assigned to " + PTXOperand::toString( type ); } - if( !PTXOperand::valid( PTXOperand::u32, b.type ) ) { + if( !PTXOperand::valid( PTXOperand::u32, b.type ) + && b.addressMode != PTXOperand::Immediate ) { return "operand 3 type " + PTXOperand::toString( b.type ) + " cannot be assigned to " + PTXOperand::toString( PTXOperand::u32 ); } - if( !PTXOperand::valid( PTXOperand::u32, b.type ) ) { + if( !PTXOperand::valid( PTXOperand::u32, c.type ) + && c.addressMode != PTXOperand::Immediate ) { return "operand 4 type " + PTXOperand::toString( c.type ) + " cannot be assigned to " + PTXOperand::toString( PTXOperand::u32 ); diff --git a/ocelot/src/parser/PTXLexer.cpp b/ocelot/src/parser/PTXLexer.cpp index 8d9c3c01..f1871cae 100644 --- a/ocelot/src/parser/PTXLexer.cpp +++ b/ocelot/src/parser/PTXLexer.cpp @@ -134,15 +134,8 @@ namespace parser CASE(TOKEN_MAXNREG) CASE(TOKEN_MAXNTID) CASE(TOKEN_MAXNCTAPERSM) - CASE(TOKEN_SM10) + CASE(TOKEN_SHADER_MODEL) CASE(TOKEN_MINNCTAPERSM) - CASE(TOKEN_SM11) - CASE(TOKEN_SM12) - CASE(TOKEN_SM13) - CASE(TOKEN_SM20) - CASE(TOKEN_SM21) - CASE(TOKEN_SM30) - CASE(TOKEN_SM35) CASE(TOKEN_MAP_F64_TO_F32) CASE(TOKEN_CONST) CASE(TOKEN_GLOBAL) diff --git a/ocelot/src/parser/PTXParser.cpp b/ocelot/src/parser/PTXParser.cpp index 636237be..0b369b70 100644 --- a/ocelot/src/parser/PTXParser.cpp +++ b/ocelot/src/parser/PTXParser.cpp @@ -376,7 +376,7 @@ namespace parser } - void PTXParser::State::ctapersm( int target, unsigned int ctas ) + void PTXParser::State::ctapersm( const char* target, unsigned int ctas ) { report( " Rule: shareModel ':' TOKEN_DECIMAL_CONSTANT" ); } @@ -598,15 +598,7 @@ namespace parser void PTXParser::State::targetElement( int token ) { report( " Rule: targetOption" ); - if( token == TOKEN_SM10 ) statement.targets.push_back( "sm_10" ); - else if( token == TOKEN_SM11 ) statement.targets.push_back( "sm_11" ); - else if( token == TOKEN_SM12 ) statement.targets.push_back( "sm_12" ); - else if( token == TOKEN_SM13 ) statement.targets.push_back( "sm_13" ); - else if( token == TOKEN_SM20 ) statement.targets.push_back( "sm_20" ); - else if( token == TOKEN_SM21 ) statement.targets.push_back( "sm_21" ); - else if( token == TOKEN_SM30 ) statement.targets.push_back( "sm_30" ); - else if( token == TOKEN_SM35 ) statement.targets.push_back( "sm_35" ); - else if( token == TOKEN_MAP_F64_TO_F32 ) + if( token == TOKEN_MAP_F64_TO_F32 ) { statement.targets.push_back( "map_f64_to_f32" ); } @@ -622,6 +614,12 @@ namespace parser } } + void PTXParser::State::targetElement( const char* target ) + { + report( " Rule: targetOption" ); + statement.targets.push_back( target ); + } + void PTXParser::State::target() { report( " Rule: TARGET targetElementList" ); diff --git a/ocelot/src/parser/ptx.ll b/ocelot/src/parser/ptx.ll index e3d747ef..927f1ccb 100644 --- a/ocelot/src/parser/ptx.ll +++ b/ocelot/src/parser/ptx.ll @@ -317,22 +317,7 @@ LABEL ({IDENTIFIER}{WHITESPACE}":") ".gl" { yylval->value = TOKEN_GL; return TOKEN_GL; } ".sys" { yylval->value = TOKEN_SYS; return TOKEN_SYS; } -"sm_10" { yylval->value = TOKEN_SM10; - return TOKEN_SM10; } -"sm_11" { yylval->value = TOKEN_SM11; - return TOKEN_SM11; } -"sm_12" { yylval->value = TOKEN_SM12; - return TOKEN_SM12; } -"sm_13" { yylval->value = TOKEN_SM13; - return TOKEN_SM13; } -"sm_20" { yylval->value = TOKEN_SM20; - return TOKEN_SM20; } -"sm_21" { yylval->value = TOKEN_SM21; - return TOKEN_SM21; } -"sm_30" { yylval->value = TOKEN_SM30; - return TOKEN_SM30; } -"sm_35" { yylval->value = TOKEN_SM35; - return TOKEN_SM35; } +"sm_"[0-9]+[a-zA-Z]* { sstrcpy( yylval->text, yytext, 1024 ); return TOKEN_SHADER_MODEL; } "map_f64_to_f32" { yylval->value = TOKEN_MAP_F64_TO_F32; return TOKEN_MAP_F64_TO_F32; } "texmode_independent" { yylval->value = TOKEN_TEXMODE_INDEPENDENT; diff --git a/ocelot/src/parser/ptxgrammar.yy b/ocelot/src/parser/ptxgrammar.yy index 32ab8e38..a5d04981 100644 --- a/ocelot/src/parser/ptxgrammar.yy +++ b/ocelot/src/parser/ptxgrammar.yy @@ -76,8 +76,8 @@ %token TOKEN_SECTION TOKEN_ADDRESS_SIZE TOKEN_WEAK %token TOKEN_MAXNREG TOKEN_MAXNTID TOKEN_MAXNCTAPERSM TOKEN_MINNCTAPERSM -%token TOKEN_SM11 TOKEN_SM12 TOKEN_SM13 TOKEN_SM20 TOKEN_MAP_F64_TO_F32 -%token TOKEN_SM21 TOKEN_SM10 TOKEN_SM30 TOKEN_SM35 +%token TOKEN_MAP_F64_TO_F32 +%token TOKEN_SHADER_MODEL %token TOKEN_TEXMODE_INDEPENDENT TOKEN_TEXMODE_UNIFIED %token TOKEN_CONST TOKEN_GLOBAL TOKEN_LOCAL TOKEN_PARAM TOKEN_PRAGMA TOKEN_PTR @@ -259,16 +259,17 @@ singleList : '{' singleListSingle '}' ',' '{' singleListSingle '}'; singleInitializer : singleList | '{' singleList '}' | '{' singleListSingle '}' | singleListSingle; -shaderModel : TOKEN_SM10 | TOKEN_SM11 | TOKEN_SM12 | TOKEN_SM13 | TOKEN_SM20 - | TOKEN_SM21 | TOKEN_SM30 | TOKEN_SM35; - floatingPointOption : TOKEN_MAP_F64_TO_F32; textureOption: TOKEN_TEXMODE_INDEPENDENT | TOKEN_TEXMODE_UNIFIED; -targetOption : shaderModel | floatingPointOption | textureOption; +targetOption : floatingPointOption | textureOption; targetElement : targetOption { state.targetElement( $1 ); +} + | TOKEN_SHADER_MODEL +{ + state.targetElement( $1 ); }; targetElementList : /* empty string */ | targetElement @@ -603,9 +604,9 @@ maxntid : TOKEN_MAXNTID TOKEN_DECIMAL_CONSTANT ',' TOKEN_DECIMAL_CONSTANT ',' state.maxntid( $2, $4, $6 ); }; -ctapersm : shaderModel ':' TOKEN_DECIMAL_CONSTANT +ctapersm : TOKEN_SHADER_MODEL ':' TOKEN_DECIMAL_CONSTANT { - state.ctapersm( $1, $3 ); + state.ctapersm( $1, $3 ); }; ctapersmList : ctapersm | ctapersmList ',' ctapersm; diff --git a/ocelot/src/parser/test/TestTargets.cpp b/ocelot/src/parser/test/TestTargets.cpp new file mode 100644 index 00000000..db31062e --- /dev/null +++ b/ocelot/src/parser/test/TestTargets.cpp @@ -0,0 +1,102 @@ +/*! \file TestTargets.cpp + \brief regression tests for .target parsing and bfi operand typing +*/ + +#include + +#include +#include +#include + +#include + +namespace test +{ + +class TestTargets: public Test +{ +public: + TestTargets() + { + name = "TestTargets"; + description = "Parses modern .target shader models, which used to be a"; + description += " lexical error, and bfi with immediate pos/len."; + } + +private: + /*! \brief parse a module, reporting the parser's own message on failure */ + bool parses(const std::string& ptx, const std::string& what) + { + std::stringstream stream(ptx); + ir::Module module; + try + { + if(!module.load(stream)) + { + status << what << ": load returned false\n"; + return false; + } + } + catch(const std::exception& e) + { + status << what << ": " << e.what() << "\n"; + return false; + } + return true; + } + + std::string kernel(const std::string& target, const std::string& body) + { + return ".version 8.0\n.target " + target + "\n.address_size 64\n" + ".visible .entry k()\n{\n" + body + "\tret;\n}\n"; + } + + /*! \brief the lexer used to stop at sm_35, so CUDA 12, which dropped + sm_35, had no target it could emit */ + bool testShaderModels() + { + const char* models[] = { "sm_10", "sm_20", "sm_35", "sm_50", "sm_61", + "sm_70", "sm_80", "sm_89", "sm_90", "sm_90a", "sm_100", "sm_120" }; + + for(auto model : models) + { + if(!parses(kernel(model, ""), model)) return false; + } + return true; + } + + /*! \brief pos and len are u32 whatever the instruction type is, but the + parser types immediates from the instruction, so .b64 made them b64 */ + bool testBfiImmediates() + { + const std::string b32 = "\t.reg .b32 %r<4>;\n" + "\tbfi.b32 %r1, %r2, %r3, 8, 16;\n"; + const std::string b64 = "\t.reg .b64 %rd<4>;\n" + "\tbfi.b64 %rd1, %rd2, %rd3, 32, 32;\n"; + + return parses(kernel("sm_50", b32), "bfi.b32") + && parses(kernel("sm_50", b64), "bfi.b64"); + } + +public: + bool doTest() + { + return testShaderModels() && testBfiImmediates(); + } +}; + +} + +int main(int argc, char** argv) +{ + hydrazine::ArgumentParser parser(argc, argv); + test::TestTargets test; + parser.description(test.testDescription()); + + parser.parse("-v", test.verbose, false, "Print out info after the test."); + parser.parse(); + + test.test(); + + return test.passed(); +}