From 014222e87977183f3e2e49a52d3a1ece4011c0a0 Mon Sep 17 00:00:00 2001 From: Jash Date: Sun, 23 Aug 2026 12:58:11 +0530 Subject: [PATCH 1/3] accept any shader model in .target The lexer enumerated sm_10 through sm_35, so every newer .target was a lexical error and ocelot aborted. That pinned callers to Kepler, which CUDA 12 no longer supports at all. Nothing reads the target string back: targetElement pushes it into statement.targets and ir/Module.cpp hardcodes "sm_21" for internally built modules. So one text-carrying token replaces the eight hardcoded ones and never needs updating for a new architecture. --- ocelot/include/ocelot/parser/PTXParser.h | 3 ++- ocelot/src/parser/PTXLexer.cpp | 9 +-------- ocelot/src/parser/PTXParser.cpp | 18 ++++++++---------- ocelot/src/parser/ptx.ll | 17 +---------------- ocelot/src/parser/ptxgrammar.yy | 17 +++++++++-------- 5 files changed, 21 insertions(+), 43 deletions(-) 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/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; From 7af6a5a7a1cd9229ec6044a58ad67f0f493daa0b Mon Sep 17 00:00:00 2001 From: Jash Date: Sun, 23 Aug 2026 12:58:11 +0530 Subject: [PATCH 2/3] bfi: pos and len are always u32 PTX types bfi's pos and len operands .u32 regardless of whether the instruction is .b32 or .b64, but the parser types immediates from the instruction type, so `bfi.b64 d, a, b, 32, 32` yielded b64 immediates and was rejected. Operand 1 already exempts immediates; operands 3 and 4 did not. Operand 4 also tested b.type instead of c.type. eval_Bfi already reads both as operandAsU32 for .b32 and .b64, so only the validator disagreed. nvcc emits bfi.b64 from sm_50 on. --- ocelot/src/ir/PTXInstruction.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) 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 ); From 55a14492a7f61c3413cb2f25d0f275b6941dc8cb Mon Sep 17 00:00:00 2001 From: Jash Date: Sun, 23 Aug 2026 19:33:07 +0530 Subject: [PATCH 3/3] test modern .target and bfi immediates regression test for the two fixes: the lexer stopped at sm_35, and bfi with immediate pos/len was rejected for .b64. --- ocelot/src/parser/test/TestTargets.cpp | 102 +++++++++++++++++++++++++ 1 file changed, 102 insertions(+) create mode 100644 ocelot/src/parser/test/TestTargets.cpp 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(); +}