accept any shader model in .target

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 );

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 );

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)

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" );

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; 

ocelot/src/parser/ptxgrammar.yy

@@ -76,8 +76,8 @@
 %token<value> TOKEN_SECTION TOKEN_ADDRESS_SIZE TOKEN_WEAK
 
 %token<value> TOKEN_MAXNREG TOKEN_MAXNTID TOKEN_MAXNCTAPERSM TOKEN_MINNCTAPERSM 
-%token<value> TOKEN_SM11 TOKEN_SM12 TOKEN_SM13 TOKEN_SM20 TOKEN_MAP_F64_TO_F32
-%token<value> TOKEN_SM21 TOKEN_SM10 TOKEN_SM30 TOKEN_SM35
+%token<value> TOKEN_MAP_F64_TO_F32
+%token<text> TOKEN_SHADER_MODEL
 %token<value> TOKEN_TEXMODE_INDEPENDENT TOKEN_TEXMODE_UNIFIED
 
 %token<value> 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( $<value>1 );
+}
+	| TOKEN_SHADER_MODEL
+{
+	state.targetElement( $<text>1 );
 };
 
 targetElementList : /* empty string */ | targetElement 
@@ -603,9 +604,9 @@ maxntid : TOKEN_MAXNTID TOKEN_DECIMAL_CONSTANT ',' TOKEN_DECIMAL_CONSTANT ','
 	state.maxntid( $<value>2, $<value>4, $<value>6 );
 };
 
-ctapersm : shaderModel ':' TOKEN_DECIMAL_CONSTANT
+ctapersm : TOKEN_SHADER_MODEL ':' TOKEN_DECIMAL_CONSTANT
 {
-	state.ctapersm( $<value>1, $<value>3 );
+	state.ctapersm( $<text>1, $<value>3 );
 };
 
 ctapersmList : ctapersm | ctapersmList ',' ctapersm;

ocelot/src/parser/test/TestTargets.cpp

@@ -0,0 +1,102 @@
+/*! \file TestTargets.cpp
+	\brief regression tests for .target parsing and bfi operand typing
+*/
+
+#include <sstream>
+
+#include <hydrazine/Test.h>
+#include <hydrazine/ArgumentParser.h>
+#include <hydrazine/Exception.h>
+
+#include <ocelot/ir/Module.h>
+
+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();
+}