support CUDA 12 sm_50 PTX

README.md

@@ -38,7 +38,7 @@ cmake ..
 make -j12
 ```
 
-Building on other distros (currently only fedora 37 was tested) requires rebuilding llvm with `-DBUILD_LLVM=ON` and turning off cuda tests with `-DBUILD_TESTS_CUDA=OFF`, as newer version of nvcc (12.0+) shipped with most distros don't support the old `sm_35` architecture.
+Building on other distros (currently only fedora 37 was tested) may require rebuilding LLVM with `-DBUILD_LLVM=ON`. CUDA tests target virtual `sm_50`, the oldest architecture supported by CUDA 12; disable them with `-DBUILD_TESTS_CUDA=OFF` when a CUDA compiler is unavailable.
 
 ### MacOS
 

ocelot/CMakeLists.txt

@@ -11,7 +11,7 @@ if (NOT APPLE AND BUILD_TESTS AND BUILD_TESTS_CUDA)
 	project(gpuocelot C CXX CUDA ASM)
 	# in order to retain the PTX code in the text executables,
 	# the target GPU architecture has to be virtual.
-	set(CMAKE_CUDA_ARCHITECTURES 35-virtual)
+	set(CMAKE_CUDA_ARCHITECTURES 50-virtual)
 	list(APPEND CUDA_NVCC_FLAGS "-Wno-deprecated-gpu-targets -Wno-deprecated-declarations")
 	set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
 else()

ocelot/src/api/test/TestEmulatedPTX.cpp

@@ -0,0 +1,119 @@
+#include <array>
+#include <cstdint>
+#include <iostream>
+#include <sstream>
+#include <string>
+
+#include <ocelot/ir/Module.h>
+#include <ocelot/parser/PTXParser.h>
+
+extern "C" void ptx_run(const char* source, int n_args, void* args[],
+	int block_x, int block_y, int block_z,
+	int grid_x, int grid_y, int grid_z, int shared_mem_size);
+
+int main()
+{
+	// Operand 4 of bfi is always u32, even when the instruction and its data operands are b64.
+	const char* invalidBfi = R"ptx(
+.version 8.7
+.target sm_50
+.address_size 64
+
+.visible .entry invalid_bfi()
+{
+	.reg .u32 %r<2>;
+	.reg .b64 %rd<4>;
+
+	mov.u32 %r1, 32;
+	mov.u64 %rd1, 1;
+	mov.u64 %rd2, 2;
+	bfi.b64 %rd3, %rd1, %rd2, %r1, %rd2;
+	ret;
+}
+)ptx";
+
+	try
+	{
+		std::stringstream source(invalidBfi);
+		ir::Module module((void*)invalidBfi, source);
+		std::cerr << "bfi.b64 accepted a b64 length operand\n";
+		return 1;
+	}
+	catch(const parser::PTXParser::Exception& error)
+	{
+		if(std::string(error.what()).find("operand 4 type b64 cannot be assigned to u32") == std::string::npos)
+		{
+			std::cerr << "bfi.b64 failed for an unexpected reason: " << error.what() << '\n';
+			return 1;
+		}
+	}
+
+	// Exercise PTX 8.7 sm_50 execution with NVRTC 12.8-style metadata and PTX-legal debug-string variants.
+	const char* ptx = R"ptx(
+.version 8.7
+.target sm_50
+.address_size 64
+
+.visible .entry insert_high_word(
+	.param .u64 insert_high_word_param_0
+)
+{
+	.reg .b32 %r<2>;
+	.reg .b64 %rd<8>;
+
+	ld.param.u64 %rd1, [insert_high_word_param_0];
+	.loc 1 7 3
+	.loc 1 4 58, function_name $L__info_string0, inlined_at 1 7 3
+	.loc 1 4 58, function_name .debug_str + 0, inlined_at 1 7 3
+	cvta.to.global.u64 %rd2, %rd1;
+	mov.u32 %r1, %tid.x;
+	mul.wide.u32 %rd3, %r1, 8;
+	add.s64 %rd4, %rd2, %rd3;
+	ld.global.u64 %rd5, [%rd4];
+	bfi.b64 %rd7, 1, %rd5, 32, 32;
+	st.global.u64 [%rd4], %rd7;
+	ret;
+}
+
+	.file 1 "insert_high_word.cu"
+	.section .debug_str
+	{
+	.b8 0
+$L__info_string0:
+	.b8 105,110,115
+	.b8 101,114,116,0
+	}
+)ptx";
+
+	{
+		std::stringstream source(ptx);
+		ir::Module module((void*)ptx, source);
+		const ir::PTXInstruction* bfi = nullptr;
+		for(const ir::PTXStatement& statement : module.statements())
+		{
+			if(statement.directive == ir::PTXStatement::Instr && statement.instruction.opcode == ir::PTXInstruction::Bfi)
+			{
+				bfi = &statement.instruction;
+				break;
+			}
+		}
+		if(bfi == nullptr || bfi->pq.type != ir::PTXOperand::b64 || bfi->a.type != ir::PTXOperand::b64 ||
+			bfi->b.type != ir::PTXOperand::u32 || bfi->c.type != ir::PTXOperand::u32)
+		{
+			std::cerr << "bfi.b64 operands were not normalized to b64, b64, u32, u32\n";
+			return 1;
+		}
+	}
+
+	std::array<std::uint64_t, 4> data{{0, 2, 4, 6}};
+	void* arguments[] = {data.data()};
+	ptx_run(ptx, 1, arguments, 4, 1, 1, 1, 1, 1, 0);
+
+	const std::array<std::uint64_t, 4> expected{{0x100000000, 0x100000002, 0x100000004, 0x100000006}};
+	if(data == expected) return 0;
+
+	std::cerr << "PTX 8.7 sm_50 emulation returned";
+	for(std::uint64_t value : data) std::cerr << ' ' << value;
+	std::cerr << "; expected high 32-bit words set to 1\n";
+	return 1;
+}

ocelot/src/ir/PTXInstruction.cpp

@@ -672,7 +672,7 @@ std::string ir::PTXInstruction::valid() const {
 					+ " cannot be assigned to " 
 					+ PTXOperand::toString( PTXOperand::u32 );
 			}
-			if( !PTXOperand::valid( PTXOperand::u32, b.type )  ) {
+			if( !PTXOperand::valid( PTXOperand::u32, c.type )  ) {
 				return "operand 4 type " + PTXOperand::toString( c.type ) 
 					+ " cannot be assigned to " 
 					+ PTXOperand::toString( PTXOperand::u32 );

ocelot/src/parser/PTXLexer.cpp

@@ -126,11 +126,14 @@ namespace parser
 			CASE(TOKEN_FILE)
 			CASE(TOKEN_VISIBLE)
 			CASE(TOKEN_LOC)
+			CASE(TOKEN_FUNCTION_NAME)
+			CASE(TOKEN_INLINED_AT)
 			CASE(TOKEN_FUNCTION)
 			CASE(TOKEN_TARGET)
 			CASE(TOKEN_VERSION)
 			CASE(TOKEN_ADDRESS_SIZE)
 			CASE(TOKEN_SECTION)
+			CASE(TOKEN_DEBUG_STR)
 			CASE(TOKEN_MAXNREG)
 			CASE(TOKEN_MAXNTID)
 			CASE(TOKEN_MAXNCTAPERSM)
@@ -143,6 +146,7 @@ namespace parser
 			CASE(TOKEN_SM21)
 			CASE(TOKEN_SM30)
 			CASE(TOKEN_SM35)
+			CASE(TOKEN_SM50)
 			CASE(TOKEN_MAP_F64_TO_F32)
 			CASE(TOKEN_CONST)
 			CASE(TOKEN_GLOBAL)

ocelot/src/parser/PTXParser.cpp

@@ -139,6 +139,27 @@ namespace parser
 	void PTXParser::State::_setImmediateTypes()
 	{
 		ir::PTXInstruction& instruction = statement.instruction;
+
+		if( instruction.opcode == ir::PTXInstruction::Bfi )
+		{
+			if( instruction.pq.addressMode == ir::PTXOperand::Immediate )
+			{
+				instruction.pq.type = instruction.type;
+			}
+			if( instruction.a.addressMode == ir::PTXOperand::Immediate )
+			{
+				instruction.a.type = instruction.type;
+			}
+			if( instruction.b.addressMode == ir::PTXOperand::Immediate )
+			{
+				instruction.b.type = ir::PTXOperand::u32;
+			}
+			if( instruction.c.addressMode == ir::PTXOperand::Immediate )
+			{
+				instruction.c.type = ir::PTXOperand::u32;
+			}
+			return;
+		}
 
 		ir::PTXOperand* sources[] =
 			{ &instruction.a, &instruction.b, &instruction.c };
@@ -606,6 +627,7 @@ namespace parser
 		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_SM50 ) statement.targets.push_back( "sm_50" );
 		else if( token == TOKEN_MAP_F64_TO_F32 )
 		{
 			statement.targets.push_back( "map_f64_to_f32" );
@@ -3000,4 +3022,3 @@ namespace parser
 }
 
 #endif
-

ocelot/src/parser/ptx.ll

@@ -268,6 +268,10 @@ LABEL ({IDENTIFIER}{WHITESPACE}":")
                                     return TOKEN_EXTERN; }
 ".file"                         { yylval->value = TOKEN_FILE; \
                                     return TOKEN_FILE; }
+"function_name"                { yylval->value = TOKEN_FUNCTION_NAME; \
+                                    return TOKEN_FUNCTION_NAME; }
+"inlined_at"                   { yylval->value = TOKEN_INLINED_AT; \
+                                    return TOKEN_INLINED_AT; }
 ".func"                         { yylval->value = TOKEN_FUNCTION; \
                                     return TOKEN_FUNCTION; }
 ".global"                       { yylval->value = TOKEN_GLOBAL; \
@@ -296,6 +300,8 @@ LABEL ({IDENTIFIER}{WHITESPACE}":")
                                     return TOKEN_SAMPLERREF; }
 ".section"                      { yylval->value = TOKEN_SECTION; \
                                     return TOKEN_SECTION; }
+".debug_str"                    { yylval->value = TOKEN_DEBUG_STR; \
+                                    return TOKEN_DEBUG_STR; }
 ".shared"                       { yylval->value = TOKEN_SHARED; \
                                     return TOKEN_SHARED;}
 ".shiftamt"                     { yylval->value = TOKEN_SHIFT_AMOUNT; \
@@ -333,6 +339,8 @@ LABEL ({IDENTIFIER}{WHITESPACE}":")
                                     return TOKEN_SM30; }
 "sm_35"                         { yylval->value = TOKEN_SM35; 
                                     return TOKEN_SM35; }
+"sm_50"                         { yylval->value = TOKEN_SM50;
+                                    return TOKEN_SM50; }
 "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

@@ -73,11 +73,12 @@
 
 %token<value> TOKEN_ENTRY TOKEN_EXTERN TOKEN_FILE TOKEN_VISIBLE TOKEN_LOC
 %token<value> TOKEN_FUNCTION TOKEN_STRUCT TOKEN_UNION TOKEN_TARGET TOKEN_VERSION
-%token<value> TOKEN_SECTION TOKEN_ADDRESS_SIZE TOKEN_WEAK
+%token<value> TOKEN_SECTION TOKEN_DEBUG_STR TOKEN_FUNCTION_NAME TOKEN_INLINED_AT
+%token<value> 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_SM21 TOKEN_SM10 TOKEN_SM30 TOKEN_SM35 TOKEN_SM50
 %token<value> TOKEN_TEXMODE_INDEPENDENT TOKEN_TEXMODE_UNIFIED
 
 %token<value> TOKEN_CONST TOKEN_GLOBAL TOKEN_LOCAL TOKEN_PARAM TOKEN_PRAGMA TOKEN_PTR
@@ -158,7 +159,7 @@ nonEntryStatement : nonEntryStatements
 };
 
 statement : initializableDeclaration | nonEntryStatement | entry | functionBody
-	| functionDeclaration;
+	| functionDeclaration | debugStringSection;
 
 statements : statement | statements statement;
 
@@ -260,7 +261,7 @@ singleInitializer : singleList |  '{' singleList '}' | '{' singleListSingle '}'
 	| singleListSingle;
 
 shaderModel : TOKEN_SM10 | TOKEN_SM11 | TOKEN_SM12 | TOKEN_SM13 | TOKEN_SM20
-	| TOKEN_SM21 | TOKEN_SM30 | TOKEN_SM35;
+	| TOKEN_SM21 | TOKEN_SM30 | TOKEN_SM35 | TOKEN_SM50;
 
 floatingPointOption : TOKEN_MAP_F64_TO_F32;
 textureOption: TOKEN_TEXMODE_INDEPENDENT | TOKEN_TEXMODE_UNIFIED;
@@ -693,11 +694,30 @@ uninitializableDeclaration : uninitializable addressableVariablePrefix
 };
 
 location : TOKEN_LOC TOKEN_DECIMAL_CONSTANT TOKEN_DECIMAL_CONSTANT 
-	TOKEN_DECIMAL_CONSTANT
+	TOKEN_DECIMAL_CONSTANT optionalInlineLocation
 {
 	state.location( $<value>2, $<value>3, $<value>4 );
 };
 
+optionalFunctionNameOffset : /* empty string */ | '+' TOKEN_DECIMAL_CONSTANT;
+
+inlineFunctionName : identifier | TOKEN_DEBUG_STR;
+
+inlineLocation : ',' TOKEN_FUNCTION_NAME inlineFunctionName optionalFunctionNameOffset
+	',' TOKEN_INLINED_AT TOKEN_DECIMAL_CONSTANT TOKEN_DECIMAL_CONSTANT
+	TOKEN_DECIMAL_CONSTANT;
+
+optionalInlineLocation : /* empty string */ | inlineLocation;
+
+debugStringBytes : TOKEN_DECIMAL_CONSTANT
+	| debugStringBytes ',' TOKEN_DECIMAL_CONSTANT;
+
+debugStringEntry : TOKEN_LABEL | TOKEN_B8 debugStringBytes;
+
+debugStringEntries : /* empty string */ | debugStringEntries debugStringEntry;
+
+debugStringSection : TOKEN_SECTION TOKEN_DEBUG_STR '{' debugStringEntries '}';
+
 label : TOKEN_LABEL optionalMetadata
 {
 	state.label( $<text>1 );