From f56f19e5aa35082413f42c201f6f5a1e32b7e5ab Mon Sep 17 00:00:00 2001 From: harish Date: Thu, 27 Aug 2026 11:26:30 +0530 Subject: [PATCH 1/4] support CUDA 12 sm_50 PTX --- README.md | 2 +- ocelot/CMakeLists.txt | 2 +- ocelot/src/api/test/TestEmulatedPTX.cpp | 58 +++++++++++++++++++++++++ ocelot/src/parser/PTXLexer.cpp | 4 ++ ocelot/src/parser/PTXParser.cpp | 1 + ocelot/src/parser/ptx.ll | 8 ++++ ocelot/src/parser/ptxgrammar.yy | 30 ++++++++++--- 7 files changed, 98 insertions(+), 7 deletions(-) create mode 100644 ocelot/src/api/test/TestEmulatedPTX.cpp diff --git a/README.md b/README.md index 6f63b870..8643ad74 100644 --- a/README.md +++ b/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 diff --git a/ocelot/CMakeLists.txt b/ocelot/CMakeLists.txt index 2662d58f..01e2e389 100644 --- a/ocelot/CMakeLists.txt +++ b/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() diff --git a/ocelot/src/api/test/TestEmulatedPTX.cpp b/ocelot/src/api/test/TestEmulatedPTX.cpp new file mode 100644 index 00000000..024a2eb4 --- /dev/null +++ b/ocelot/src/api/test/TestEmulatedPTX.cpp @@ -0,0 +1,58 @@ +#include +#include + +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() +{ + // 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 increment( + .param .u64 increment_param_0 +) +{ + .reg .b32 %r<4>; + .reg .b64 %rd<5>; + + ld.param.u64 %rd1, [increment_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, 4; + add.s64 %rd4, %rd2, %rd3; + ld.global.u32 %r2, [%rd4]; + add.s32 %r3, %r2, 1; + st.global.u32 [%rd4], %r3; + ret; +} + + .file 1 "increment.cu" + .section .debug_str + { + .b8 0 +$L__info_string0: + .b8 105,110,99,114,101 + .b8 109,101,110,116,0 + } +)ptx"; + + std::array data{{1, 2, 3, 4}}; + void* arguments[] = {data.data()}; + ptx_run(ptx, 1, arguments, 4, 1, 1, 1, 1, 1, 0); + + const std::array expected{{2, 3, 4, 5}}; + if(data == expected) return 0; + + std::cerr << "PTX 8.7 sm_50 emulation returned"; + for(unsigned int value : data) std::cerr << ' ' << value; + std::cerr << "; expected 2 3 4 5\n"; + return 1; +} diff --git a/ocelot/src/parser/PTXLexer.cpp b/ocelot/src/parser/PTXLexer.cpp index 8d9c3c01..ff7ddd71 100644 --- a/ocelot/src/parser/PTXLexer.cpp +++ b/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) diff --git a/ocelot/src/parser/PTXParser.cpp b/ocelot/src/parser/PTXParser.cpp index 636237be..7f066676 100644 --- a/ocelot/src/parser/PTXParser.cpp +++ b/ocelot/src/parser/PTXParser.cpp @@ -606,6 +606,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" ); diff --git a/ocelot/src/parser/ptx.ll b/ocelot/src/parser/ptx.ll index e3d747ef..00e06c8f 100644 --- a/ocelot/src/parser/ptx.ll +++ b/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; diff --git a/ocelot/src/parser/ptxgrammar.yy b/ocelot/src/parser/ptxgrammar.yy index 32ab8e38..5dd563ab 100644 --- a/ocelot/src/parser/ptxgrammar.yy +++ b/ocelot/src/parser/ptxgrammar.yy @@ -73,11 +73,12 @@ %token TOKEN_ENTRY TOKEN_EXTERN TOKEN_FILE TOKEN_VISIBLE TOKEN_LOC %token TOKEN_FUNCTION TOKEN_STRUCT TOKEN_UNION TOKEN_TARGET TOKEN_VERSION -%token TOKEN_SECTION TOKEN_ADDRESS_SIZE TOKEN_WEAK +%token TOKEN_SECTION TOKEN_DEBUG_STR TOKEN_FUNCTION_NAME TOKEN_INLINED_AT +%token 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_SM21 TOKEN_SM10 TOKEN_SM30 TOKEN_SM35 TOKEN_SM50 %token TOKEN_TEXMODE_INDEPENDENT TOKEN_TEXMODE_UNIFIED %token 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( $2, $3, $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( $1 ); From 562d4e60c4fa8e5c53512bcc2088a595f78844db Mon Sep 17 00:00:00 2001 From: harish Date: Thu, 27 Aug 2026 11:48:29 +0530 Subject: [PATCH 2/4] fix BFI immediate operand types --- ocelot/src/api/test/TestEmulatedPTX.cpp | 56 ++++++++++++++++++------- ocelot/src/ir/PTXInstruction.cpp | 2 +- ocelot/src/parser/PTXParser.cpp | 24 ++++++++++- 3 files changed, 64 insertions(+), 18 deletions(-) diff --git a/ocelot/src/api/test/TestEmulatedPTX.cpp b/ocelot/src/api/test/TestEmulatedPTX.cpp index 024a2eb4..11aa3ad7 100644 --- a/ocelot/src/api/test/TestEmulatedPTX.cpp +++ b/ocelot/src/api/test/TestEmulatedPTX.cpp @@ -1,5 +1,9 @@ #include +#include #include +#include + +#include extern "C" void ptx_run(const char* source, int n_args, void* args[], int block_x, int block_y, int block_z, @@ -13,46 +17,66 @@ int main() .target sm_50 .address_size 64 -.visible .entry increment( - .param .u64 increment_param_0 +.visible .entry insert_high_word( + .param .u64 insert_high_word_param_0 ) { - .reg .b32 %r<4>; - .reg .b64 %rd<5>; + .reg .b32 %r<2>; + .reg .b64 %rd<8>; - ld.param.u64 %rd1, [increment_param_0]; + 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, 4; + mul.wide.u32 %rd3, %r1, 8; add.s64 %rd4, %rd2, %rd3; - ld.global.u32 %r2, [%rd4]; - add.s32 %r3, %r2, 1; - st.global.u32 [%rd4], %r3; + ld.global.u64 %rd5, [%rd4]; + bfi.b64 %rd7, 1, %rd5, 32, 32; + st.global.u64 [%rd4], %rd7; ret; } - .file 1 "increment.cu" + .file 1 "insert_high_word.cu" .section .debug_str { .b8 0 $L__info_string0: - .b8 105,110,99,114,101 - .b8 109,101,110,116,0 + .b8 105,110,115 + .b8 101,114,116,0 } )ptx"; - std::array data{{1, 2, 3, 4}}; + { + 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 data{{0, 2, 4, 6}}; void* arguments[] = {data.data()}; ptx_run(ptx, 1, arguments, 4, 1, 1, 1, 1, 1, 0); - const std::array expected{{2, 3, 4, 5}}; + const std::array expected{{0x100000000, 0x100000002, 0x100000004, 0x100000006}}; if(data == expected) return 0; std::cerr << "PTX 8.7 sm_50 emulation returned"; - for(unsigned int value : data) std::cerr << ' ' << value; - std::cerr << "; expected 2 3 4 5\n"; + for(std::uint64_t value : data) std::cerr << ' ' << value; + std::cerr << "; expected high 32-bit words set to 1\n"; return 1; } diff --git a/ocelot/src/ir/PTXInstruction.cpp b/ocelot/src/ir/PTXInstruction.cpp index 099a0493..5df5118d 100644 --- a/ocelot/src/ir/PTXInstruction.cpp +++ b/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 ); diff --git a/ocelot/src/parser/PTXParser.cpp b/ocelot/src/parser/PTXParser.cpp index 7f066676..ed71fada 100644 --- a/ocelot/src/parser/PTXParser.cpp +++ b/ocelot/src/parser/PTXParser.cpp @@ -32,6 +32,8 @@ #undef REPORT_BASE #endif + + #define REPORT_BASE 0 /*! \brief A namespace for parsing PTX */ @@ -139,6 +141,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 }; @@ -3001,4 +3024,3 @@ namespace parser } #endif - From 52a290cdf3906ab3b4d3a0a1b596177efc38e5a4 Mon Sep 17 00:00:00 2001 From: harish Date: Thu, 27 Aug 2026 11:54:21 +0530 Subject: [PATCH 3/4] remove stray parser whitespace --- ocelot/src/parser/PTXParser.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/ocelot/src/parser/PTXParser.cpp b/ocelot/src/parser/PTXParser.cpp index ed71fada..e8d68bd1 100644 --- a/ocelot/src/parser/PTXParser.cpp +++ b/ocelot/src/parser/PTXParser.cpp @@ -32,8 +32,6 @@ #undef REPORT_BASE #endif - - #define REPORT_BASE 0 /*! \brief A namespace for parsing PTX */ From 664c83145bf0f38ec7f2a635eef152c3632f1c9c Mon Sep 17 00:00:00 2001 From: harish Date: Thu, 27 Aug 2026 12:13:06 +0530 Subject: [PATCH 4/4] test invalid BFI length operand --- ocelot/src/api/test/TestEmulatedPTX.cpp | 37 +++++++++++++++++++++++++ 1 file changed, 37 insertions(+) diff --git a/ocelot/src/api/test/TestEmulatedPTX.cpp b/ocelot/src/api/test/TestEmulatedPTX.cpp index 11aa3ad7..06ec5a92 100644 --- a/ocelot/src/api/test/TestEmulatedPTX.cpp +++ b/ocelot/src/api/test/TestEmulatedPTX.cpp @@ -2,8 +2,10 @@ #include #include #include +#include #include +#include extern "C" void ptx_run(const char* source, int n_args, void* args[], int block_x, int block_y, int block_z, @@ -11,6 +13,41 @@ extern "C" void ptx_run(const char* source, int n_args, void* args[], 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