From 014222e87977183f3e2e49a52d3a1ece4011c0a0 Mon Sep 17 00:00:00 2001 From: Jash Date: Sun, 23 Aug 2026 12:58:11 +0530 Subject: [PATCH 1/5] 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/5] 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/5] 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(); +} From 5b436dd2fe5c6bcb887dcee3ac548a85b4fe0fd9 Mon Sep 17 00:00:00 2001 From: Jash Date: Sun, 23 Aug 2026 19:53:58 +0530 Subject: [PATCH 4/5] implement f16 cvt to and from f16 threw "conversion not implemented", and add, sub, mul, abs, neg, min, max, fma, set, setp, mov and ex2 rejected f16 as an instruction type, so nothing targeting sm_53 or newer could use half. For a CUDA 12 toolchain that is every target, since sm_35 is gone. f32 represents every f16 exactly, so operandAsHalf widens an operand and the existing f32 paths are reused: cvt takes f16 as a source by falling through to the f32 case, and set and setp likewise. Destinations narrow with floatToHalf. The exact result of an f16 add or mul also fits in f32 (11 bit significands, so a product needs 22 of the 24 available), so computing in f32 and rounding once to f16 is correctly rounded rather than double rounded. ex2 is .approx in PTX, so the f32 approximation is in spec. Reading an operand cannot key off the declared type. tinygrad's PTX renderer declares halves .f16; nvrtc declares them .b16 and puts the op in inline asm. The half is in the low 16 bits either way. Immediates keep the f32 handling: the parser types an immediate from the instruction and the lexer has no 0H half literal, so a constant reaching an .f16 instruction was written as a float and lives in imm_single, and reading it as bits picks up the wrong half of the union. cvt needed one more thing: relaxedConvert stores the mnemonic's type in a.relaxedType and leaves a.type as the register declaration, so operandAsF32 has to honour either. Without that, half loads on the nvrtc path came back as denormals like 2.15e-41, which is 0x3C00 sitting in a float. set also accepts a .b16 destination, since that is the register nvrtc writes the 1.0h/0.0h into. halfToFloat and floatToHalf were checked against numpy: all 65536 half bit patterns widen to the exact f32 bit pattern, and 531082 narrowing cases including every exact tie midpoint round identically, ties to even. mad keeps rejecting f16, which is correct: PTX has no mad.f16, half fma is spelled fma.rn.f16 and that is implemented. ocelot has no tanh opcode for any type, so tanh.approx.f16 still does not parse, and half atomics are unimplemented and abort. Neither is reached by tinygrad or by the parts of cuda_fp16.hpp the tests exercise. tinygrad's emulated CI, test/backend, both emulator consumers: ptx cell nv cell master 1738 pass 453 skip 1669 pass 458 skip sm_35 1738 pass 453 skip 1669 pass 458 skip sm_53 1738 pass 443 skip 1669 pass 449 skip Zero failures everywhere. sm_35 is unchanged on both cells, so this is a no-op at the current target. The freed skips at sm_53 are the fp16 tests, which have never been able to run: tinygrad disables half below sm_53, so a Kepler-reporting mock meant the NV emulator never tested fp16 at all. The nv cell at sm_53 runs on nvrtc 12.9 with matching cudart headers. --- .../ocelot/executive/CooperativeThreadArray.h | 1 + .../src/executive/CooperativeThreadArray.cpp | 267 +++++++++++++++++- ocelot/src/ir/PTXInstruction.cpp | 48 ++-- 3 files changed, 276 insertions(+), 40 deletions(-) diff --git a/ocelot/include/ocelot/executive/CooperativeThreadArray.h b/ocelot/include/ocelot/executive/CooperativeThreadArray.h index 84e51a00..4046d2d3 100644 --- a/ocelot/include/ocelot/executive/CooperativeThreadArray.h +++ b/ocelot/include/ocelot/executive/CooperativeThreadArray.h @@ -425,6 +425,7 @@ namespace executive { ir::PTXS64 operandAsS64(int, const ir::PTXOperand &); ir::PTXF32 operandAsF32(int, const ir::PTXOperand &); + ir::PTXF32 operandAsHalf(int, const ir::PTXOperand &); ir::PTXF64 operandAsF64(int, const ir::PTXOperand &); ir::PTXB8 operandAsB8(int, const ir::PTXOperand &); diff --git a/ocelot/src/executive/CooperativeThreadArray.cpp b/ocelot/src/executive/CooperativeThreadArray.cpp index 81acd6dd..4246860f 100644 --- a/ocelot/src/executive/CooperativeThreadArray.cpp +++ b/ocelot/src/executive/CooperativeThreadArray.cpp @@ -859,6 +859,52 @@ ir::PTXS64 executive::CooperativeThreadArray::getRegAsS64(int threadID, \param threadID ID of the active thread \reg register index */ +/*! + PTX keeps .f16 values as an IEEE754 binary16 bit pattern in the low half of the + register slot. f32 represents every f16 exactly, so widening is lossless and a + cvt from f16 is exactly the cvt from the widened f32. +*/ +static ir::PTXF32 halfToFloat(ir::PTXU16 h) { + ir::PTXU32 sign = (ir::PTXU32)(h & 0x8000) << 16; + ir::PTXU32 exp = (h >> 10) & 0x1f, mant = h & 0x3ff, bits; + if (exp == 0) { + if (mant == 0) bits = sign; + else { + exp = 127 - 15 + 1; + while ((mant & 0x400) == 0) { mant <<= 1; --exp; } + bits = sign | (exp << 23) | ((mant & 0x3ff) << 13); + } + } + else if (exp == 0x1f) bits = sign | 0x7f800000 | (mant << 13); + else bits = sign | ((exp - 15 + 127) << 23) | (mant << 13); + return hydrazine::bit_cast(bits); +} + +/*! binary32 to binary16, round to nearest even, which is the PTX cvt.rn default. */ +static ir::PTXU16 floatToHalf(ir::PTXF32 f) { + ir::PTXU32 bits = hydrazine::bit_cast(f); + ir::PTXU16 sign = (ir::PTXU16)((bits >> 16) & 0x8000); + ir::PTXU32 rawexp = (bits >> 23) & 0xff, mant = bits & 0x7fffff; + if (rawexp == 0xff) { + return sign | 0x7c00 | (mant ? (ir::PTXU16)((mant >> 13) | 0x200) : 0); + } + int exp = (int)rawexp - 127 + 15; + if (exp >= 0x1f) return sign | 0x7c00; + if (exp <= 0) { + if (exp < -10) return sign; + mant |= 0x800000; + int shift = 14 - exp; + ir::PTXU32 h = mant >> shift; + ir::PTXU32 rem = mant & ((1u << shift) - 1), half = 1u << (shift - 1); + if (rem > half || (rem == half && (h & 1))) ++h; + return sign | (ir::PTXU16)h; + } + ir::PTXU16 h = (ir::PTXU16)((exp << 10) | (mant >> 13)); + ir::PTXU32 rem = mant & 0x1fff; + if (rem > 0x1000 || (rem == 0x1000 && (h & 1))) ++h; + return sign | h; +} + ir::PTXF32 executive::CooperativeThreadArray::getRegAsF32(int threadID, ir::PTXOperand::RegisterType reg) { ir::PTXF32 r = *( (ir::PTXF32*)( @@ -1496,10 +1542,35 @@ ir::PTXS64 executive::CooperativeThreadArray::operandAsS64(int threadID, return 0; } +/*! + Read an operand of an .f16 instruction as f32. + + The register may be declared .f16, which is what tinygrad's PTX renderer emits, or + .b16, which is what nvrtc emits when it puts the op in inline asm. The half sits in + the low 16 bits either way, so go through the bits rather than the declared type. + + Immediates keep the normal f32 handling: the parser types an immediate from the + instruction, and the lexer has no 0H half literal, so a constant reaching an .f16 + instruction was written as a float and lives in imm_single. Reading it as bits would + pick up the wrong half of the union. +*/ +ir::PTXF32 executive::CooperativeThreadArray::operandAsHalf(int threadID, + const ir::PTXOperand &op) { + if (op.addressMode == ir::PTXOperand::Immediate) return operandAsF32(threadID, op); + return halfToFloat(operandAsB16(threadID, op)); +} + ir::PTXF32 executive::CooperativeThreadArray::operandAsF32(int threadID, const ir::PTXOperand &op) { switch (op.addressMode) { case ir::PTXOperand::Register: + // relaxedType carries the type from the mnemonic when it differs from the + // register declaration, which is how nvrtc's cvt.f32.f16 on a .b16 register + // arrives. Either one saying f16 means the bits are a half. + if (op.type == ir::PTXOperand::f16 + || op.relaxedType == ir::PTXOperand::f16) { + return halfToFloat(getRegAsB16(threadID, op.reg)); + } return getRegAsF32(threadID, op.reg); case ir::PTXOperand::Immediate: return (ir::PTXF32)(op.imm_single); @@ -1657,7 +1728,14 @@ void executive::CooperativeThreadArray::setFunctionParameter(int threadID, void executive::CooperativeThreadArray::eval_Abs(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + setRegAsB16(threadID, instr.d.reg, + floatToHalf(CTAAbs(operandAsHalf(threadID, instr.a)))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; @@ -1717,7 +1795,16 @@ void executive::CooperativeThreadArray::eval_Abs(CTAContext &context, void executive::CooperativeThreadArray::eval_Add(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + ir::PTXF32 a = operandAsHalf(threadID, instr.a), + b = operandAsHalf(threadID, instr.b); + setRegAsB16(threadID, instr.d.reg, + floatToHalf(sat(instr.modifier, a + b))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; ir::PTXF32 d, a = ftz(instr.modifier, operandAsF32(threadID, instr.a)), @@ -3009,6 +3096,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, d); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsB8(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3059,6 +3153,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsU64(threadID, instr.d.reg, a); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsS8(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3126,6 +3227,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, d); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsB16(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3192,6 +3300,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsU64(threadID, instr.d.reg, a); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsS16(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3274,6 +3389,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, d); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsU32(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3354,6 +3476,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, a); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsS32(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3449,6 +3578,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, a); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsS64(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3546,6 +3682,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, d); } break; + case ir::PTXOperand::f16: + { + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(operandAsU64(threadID, instr.a), + instr.modifier))); + } + break; case ir::PTXOperand::f32: { setRegAsF32(threadID, instr.d.reg, @@ -3567,6 +3710,7 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, } } break; + case ir::PTXOperand::f16: // fall through, widened by operandAsF32 case ir::PTXOperand::f32: { switch (instr.type) { @@ -3727,6 +3871,17 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, d); } break; + case ir::PTXOperand::f16: + { + ir::PTXF32 a = operandAsF32(threadID, instr.a); + + a = roundToInt(a, instr.modifier, context, + instr); + + setRegAsB16(threadID, instr.d.reg, + floatToHalf(sat(instr.modifier, a))); + } + break; case ir::PTXOperand::f32: { ir::PTXF32 a = operandAsF32(threadID, instr.a); @@ -3914,6 +4069,13 @@ void executive::CooperativeThreadArray::eval_Cvt(CTAContext &context, setRegAsS64(threadID, instr.d.reg, d); } break; + case ir::PTXOperand::f16: + { + ir::PTXF64 a = operandAsF64(threadID, instr.a); + setRegAsB16(threadID, instr.d.reg, + floatToHalf(toF32(a, instr.modifier))); + } + break; case ir::PTXOperand::f32: { ir::PTXF64 a = operandAsF64(threadID, instr.a); @@ -4373,7 +4535,15 @@ void executive::CooperativeThreadArray::eval_Div(CTAContext &context, void executive::CooperativeThreadArray::eval_Ex2(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + ir::PTXF32 a = operandAsHalf(threadID, instr.a); + setRegAsB16(threadID, instr.d.reg, + floatToHalf(hydrazine::exp2f(a))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; @@ -4403,7 +4573,17 @@ void executive::CooperativeThreadArray::eval_Exit(CTAContext &context, void executive::CooperativeThreadArray::eval_Fma(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int tid = 0; tid < threadCount; tid++) { + if (!context.predicated(tid, instr)) continue; + ir::PTXF32 a = operandAsHalf(tid, instr.a), + b = operandAsHalf(tid, instr.b), + c = operandAsHalf(tid, instr.c); + setRegAsB16(tid, instr.d.reg, + floatToHalf(sat(instr.modifier, a * b + c))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int tid = 0; tid < threadCount; tid++) { if (!context.predicated(tid, instr)) continue; ir::PTXF32 d = 0, @@ -5316,7 +5496,19 @@ void executive::CooperativeThreadArray::eval_Mad(CTAContext &context, void executive::CooperativeThreadArray::eval_Max(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + ir::PTXF32 a = operandAsHalf(threadID, instr.a), + b = operandAsHalf(threadID, instr.b); + ir::PTXF32 d; + if (hydrazine::isnan(a)) d = b; + else if (hydrazine::isnan(b)) d = a; + else d = (a > b) ? a : b; + setRegAsB16(threadID, instr.d.reg, floatToHalf(d)); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; @@ -5433,7 +5625,18 @@ void executive::CooperativeThreadArray::eval_Max(CTAContext &context, void executive::CooperativeThreadArray::eval_Min(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + ir::PTXF32 d, a = operandAsHalf(threadID, instr.a), + b = operandAsHalf(threadID, instr.b); + if (hydrazine::isnan(a)) d = b; + else if (hydrazine::isnan(b)) d = a; + else d = (a < b) ? a : b; + setRegAsB16(threadID, instr.d.reg, floatToHalf(d)); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; @@ -5808,6 +6011,7 @@ void executive::CooperativeThreadArray::eval_Mov_imm(CTAContext &context, case ir::PTXOperand::u16: case ir::PTXOperand::s16: case ir::PTXOperand::b16: + case ir::PTXOperand::f16: // PTX materializes a half constant as mov.b16 into an .f16 reg { ir::PTXU16 a = operandAsU16(threadID, instr.a); setRegAsU16(threadID, instr.d.reg, a); @@ -5950,7 +6154,16 @@ void executive::CooperativeThreadArray::eval_Mul24(CTAContext &context, const ir */ void executive::CooperativeThreadArray::eval_Mul(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + ir::PTXF32 a = operandAsHalf(threadID, instr.a), + b = operandAsHalf(threadID, instr.b); + setRegAsB16(threadID, instr.d.reg, + floatToHalf(sat(instr.modifier, a * b))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; @@ -6113,7 +6326,14 @@ void executive::CooperativeThreadArray::eval_Mul(CTAContext &context, const ir:: */ void executive::CooperativeThreadArray::eval_Neg(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + setRegAsB16(threadID, instr.d.reg, + floatToHalf(-operandAsHalf(threadID, instr.a))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; @@ -7132,14 +7352,18 @@ void executive::CooperativeThreadArray::eval_SetP(CTAContext &context, } break; + case ir::PTXOperand::f16: // fall through, read as half below // single-precision float case ir::PTXOperand::f32: { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; - ir::PTXF32 a = ftz(instr.modifier, operandAsF32(threadID, instr.a)), - b = ftz(instr.modifier, operandAsF32(threadID, instr.b)); + bool half = instr.type == ir::PTXOperand::f16; + ir::PTXF32 a = half ? operandAsHalf(threadID, instr.a) + : ftz(instr.modifier, operandAsF32(threadID, instr.a)), + b = half ? operandAsHalf(threadID, instr.b) + : ftz(instr.modifier, operandAsF32(threadID, instr.b)); bool c = true; // read operator somehow bool t = false; @@ -7346,7 +7570,12 @@ void executive::CooperativeThreadArray::eval_Set(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - switch (instr.a.type) { + // set.CmpOp.dtype.f16 compares halves, but nvrtc declares the registers .b16 and + // the mnemonic type lands in relaxedType, so switch on that when it says f16. + ir::PTXOperand::DataType sourceType = instr.a.type; + if (instr.a.relaxedType == ir::PTXOperand::f16) sourceType = ir::PTXOperand::f16; + + switch (sourceType) { // unsigned int types [extended to 64-bit uint] case ir::PTXOperand::b16: @@ -7551,6 +7780,7 @@ void executive::CooperativeThreadArray::eval_Set(CTAContext &context, break; // single-precision float + case ir::PTXOperand::f16: // fall through, operandAsF32 widens case ir::PTXOperand::f32: { for (int threadID = 0; threadID < threadCount; threadID++) { @@ -7647,6 +7877,10 @@ void executive::CooperativeThreadArray::eval_Set(CTAContext &context, case ir::PTXOperand::u32: setRegAsU32(threadID, instr.d.reg, (t ? 0xFFFFFFFF : 0x00)); break; + case ir::PTXOperand::f16: + setRegAsB16(threadID, instr.d.reg, + floatToHalf(t ? 1.0f : 0.0f)); + break; case ir::PTXOperand::f32: setRegAsF32(threadID, instr.d.reg, (t ? 1.0f : 0.0f)); break; @@ -8750,7 +8984,16 @@ void executive::CooperativeThreadArray::eval_St(CTAContext &context, void executive::CooperativeThreadArray::eval_Sub(CTAContext &context, const ir::PTXInstruction &instr) { trace(); - if (instr.type == ir::PTXOperand::f32) { + if (instr.type == ir::PTXOperand::f16) { + for (int threadID = 0; threadID < threadCount; threadID++) { + if (!context.predicated(threadID, instr)) continue; + ir::PTXF32 a = operandAsHalf(threadID, instr.a), + b = operandAsHalf(threadID, instr.b); + setRegAsB16(threadID, instr.d.reg, + floatToHalf(sat(instr.modifier, a - b))); + } + } + else if (instr.type == ir::PTXOperand::f32) { for (int threadID = 0; threadID < threadCount; threadID++) { if (!context.predicated(threadID, instr)) continue; diff --git a/ocelot/src/ir/PTXInstruction.cpp b/ocelot/src/ir/PTXInstruction.cpp index c42c7506..7982801d 100644 --- a/ocelot/src/ir/PTXInstruction.cpp +++ b/ocelot/src/ir/PTXInstruction.cpp @@ -481,8 +481,8 @@ std::string ir::PTXInstruction::valid() const { switch (opcode) { case Abs: { if ( !( type == PTXOperand::s16 || type == PTXOperand::s32 || - type == PTXOperand::s64 || type == PTXOperand::f32 || - type == PTXOperand::f64 ) ) { + type == PTXOperand::s64 || type == PTXOperand::f16 || + type == PTXOperand::f32 || type == PTXOperand::f64 ) ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -504,8 +504,7 @@ std::string ir::PTXInstruction::valid() const { } case Add: { if ( !( type != PTXOperand::s8 && type != PTXOperand::u8 && - type != PTXOperand::b8 && type != PTXOperand::f16 - && type != PTXOperand::pred ) ) { + type != PTXOperand::b8 && type != PTXOperand::pred ) ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -923,7 +922,7 @@ std::string ir::PTXInstruction::valid() const { break; } case Ex2: { - if( !( type == PTXOperand::f32 ) ) { + if( !( type == PTXOperand::f32 || type == PTXOperand::f16 ) ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -947,7 +946,8 @@ std::string ir::PTXInstruction::valid() const { break; } case Fma: { - if (!(type == ir::PTXOperand::f32 || type == ir::PTXOperand::f64)) { + if (!(type == ir::PTXOperand::f16 || type == ir::PTXOperand::f32 + || type == ir::PTXOperand::f64)) { return "invalid instruction type " + PTXOperand::toString( type ); } if( !PTXOperand::valid( type, d.type ) ) { @@ -1124,8 +1124,7 @@ std::string ir::PTXInstruction::valid() const { } case Max: { if( !( type != PTXOperand::s8 && type != PTXOperand::u8 && - type != PTXOperand::b8 && type != PTXOperand::f16 - && type != PTXOperand::pred ) ) { + type != PTXOperand::b8 && type != PTXOperand::pred ) ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -1154,8 +1153,7 @@ std::string ir::PTXInstruction::valid() const { } case Min: { if( !( type != PTXOperand::s8 && type != PTXOperand::u8 && - type != PTXOperand::b8 && type != PTXOperand::f16 - && type != PTXOperand::pred ) ) { + type != PTXOperand::b8 && type != PTXOperand::pred ) ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -1180,14 +1178,8 @@ std::string ir::PTXInstruction::valid() const { break; } case Mov: { - if ( ( a.type == PTXOperand::f16 ) && - a.addressMode != PTXOperand::Address && - a.addressMode != PTXOperand::Immediate ) { - return "invalid type for operand A " - + PTXOperand::toString( a.type ); - } if ( !( d.type != PTXOperand::s8 && d.type != PTXOperand::u8 - && d.type != PTXOperand::b8 && d.type != PTXOperand::f16 ) ) { + && d.type != PTXOperand::b8 ) ) { return "invalid type for operand D " + PTXOperand::toString( d.type ); } @@ -1224,8 +1216,7 @@ std::string ir::PTXInstruction::valid() const { } case Mul: { if( type == PTXOperand::s8 || type == PTXOperand::u8 - || type == PTXOperand::b8 || type == PTXOperand::f16 - || type == PTXOperand::pred ) { + || type == PTXOperand::b8 || type == PTXOperand::pred ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -1265,8 +1256,8 @@ std::string ir::PTXInstruction::valid() const { } case Neg: { if( type != PTXOperand::s16 && type != PTXOperand::s32 && - type != PTXOperand::s64 && type != PTXOperand::f32 && - type != PTXOperand::f64 ) { + type != PTXOperand::s64 && type != PTXOperand::f16 && + type != PTXOperand::f32 && type != PTXOperand::f64 ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -1577,15 +1568,17 @@ std::string ir::PTXInstruction::valid() const { && type != PTXOperand::s64 && type != PTXOperand::u16 && type != PTXOperand::u32 && type != PTXOperand::u64 && type != PTXOperand::b16 && type != PTXOperand::b32 - && type != PTXOperand::b64 && type != PTXOperand::f32 - && type != PTXOperand::f64 ) { + && type != PTXOperand::b64 && type != PTXOperand::f16 + && type != PTXOperand::f32 && type != PTXOperand::f64 ) { return "invalid instruction type " + PTXOperand::toString( type ); } + // an .f16 set writes a half, and nvrtc declares that register .b16 if( d.type != PTXOperand::s32 && d.type != PTXOperand::f32 + && d.type != PTXOperand::f16 && d.type != PTXOperand::b16 && d.type != PTXOperand::u32 ) { return "operand D type " + PTXOperand::toString( d.type ) - + " invalid (must be u32, s32, or f32)"; + + " invalid (must be u32, s32, f16, or f32)"; } if( c.type != PTXOperand::pred && c.addressMode != PTXOperand::Invalid ) { @@ -1616,8 +1609,8 @@ std::string ir::PTXInstruction::valid() const { && type != PTXOperand::s64 && type != PTXOperand::u16 && type != PTXOperand::u32 && type != PTXOperand::u64 && type != PTXOperand::b16 && type != PTXOperand::b32 - && type != PTXOperand::b64 && type != PTXOperand::f32 - && type != PTXOperand::f64 ) { + && type != PTXOperand::b64 && type != PTXOperand::f16 + && type != PTXOperand::f32 && type != PTXOperand::f64 ) { return "invalid instruction type " + PTXOperand::toString( type ); } @@ -1864,8 +1857,7 @@ std::string ir::PTXInstruction::valid() const { } case Sub: { if ( !( type != PTXOperand::s8 && type != PTXOperand::u8 && - type != PTXOperand::b8 && type != PTXOperand::f16 - && type != PTXOperand::pred ) ) { + type != PTXOperand::b8 && type != PTXOperand::pred ) ) { return "invalid instruction type " + PTXOperand::toString( type ); } From 2f13ecba8601e4616e9959c051eaa0af77902a78 Mon Sep 17 00:00:00 2001 From: Jash Date: Sun, 23 Aug 2026 19:53:58 +0530 Subject: [PATCH 5/5] test f16 covers cvt both directions including the .b16 declaration nvrtc emits, ties to even on narrowing, the arithmetic and compare ops, and the mov.b16 immediate that silently did nothing. --- .../src/executive/test/TestInstructions.cpp | 159 ++++++++++++++++++ 1 file changed, 159 insertions(+) diff --git a/ocelot/src/executive/test/TestInstructions.cpp b/ocelot/src/executive/test/TestInstructions.cpp index ce779bb8..cc136a53 100644 --- a/ocelot/src/executive/test/TestInstructions.cpp +++ b/ocelot/src/executive/test/TestInstructions.cpp @@ -209,6 +209,164 @@ class TestInstructions: public Test { /*! Tests several forms of the abs instruction */ + /*! \brief half bit patterns used below */ + enum { H0 = 0x0000, H1 = 0x3C00, H2 = 0x4000, H3 = 0x4200, H4 = 0x4400, + HNEG1 = 0xBC00 }; + + /*! \brief an f16 operand. tinygrad's PTX renderer declares halves .f16; + nvrtc declares them .b16 and leaves the type in relaxedType. */ + PTXOperand f16reg(PTXOperand::RegisterType r, bool declaredB16 = false) { + PTXOperand op = reg("h", declaredB16 ? PTXOperand::b16 + : PTXOperand::f16, r); + if (declaredB16) op.relaxedType = PTXOperand::f16; + return op; + } + + bool checkHalf(const char* what, PTXOperand::RegisterType r, PTXU16 want) { + for (int t = 0; t < threadCount; t++) { + PTXU16 got = cta->getRegAsB16(t, r); + if (got != want) { + status << what << " failed (thread " << t << "): expected 0x" + << std::hex << want << ", got 0x" << got << std::dec << "\n"; + return false; + } + } + return true; + } + + void setHalf(PTXOperand::RegisterType r, PTXU16 bits) { + for (int t = 0; t < threadCount; t++) cta->setRegAsB16(t, r, bits); + } + + /*! \brief cvt and the ALU rejected or mishandled f16, so no target from + sm_53 on could use half at all */ + bool test_f16() { + PTXInstruction ins; + cta->reset(); + + // cvt.f32.f16, with .f16 and with the .b16 declaration nvrtc emits + for (int b16 = 0; b16 < 2; b16++) { + ins = PTXInstruction(); + ins.opcode = PTXInstruction::Cvt; + ins.type = PTXOperand::f32; + ins.d = reg("f", PTXOperand::f32, 0); + ins.a = f16reg(1, b16 != 0); + setHalf(1, H2); + cta->eval_Cvt(cta->getActiveContext(), ins); + for (int t = 0; t < threadCount; t++) { + if (cta->getRegAsF32(t, 0) != 2.0f) { + status << "cvt.f32.f16 failed for " + << (b16 ? ".b16" : ".f16") << " operand: got " + << cta->getRegAsF32(t, 0) << "\n"; + return false; + } + } + } + + // cvt.rn.f16.f32 rounds to nearest even, so the two midpoints go to the + // even mantissa rather than both going up + const float mids[2] = { 1.00048828125f, 1.00146484375f }; + const PTXU16 wants[2] = { 0x3C00, 0x3C02 }; + for (int i = 0; i < 2; i++) { + ins = PTXInstruction(); + ins.opcode = PTXInstruction::Cvt; + ins.type = PTXOperand::f16; + ins.modifier = PTXInstruction::rn; + ins.d = f16reg(0); + ins.a = reg("f", PTXOperand::f32, 1); + for (int t = 0; t < threadCount; t++) cta->setRegAsF32(t, 1, mids[i]); + cta->eval_Cvt(cta->getActiveContext(), ins); + if (!checkHalf("cvt.rn.f16.f32 ties to even", 0, wants[i])) return false; + } + + // binary arithmetic, 2 op 1, with mixed operand declarations + ins = PTXInstruction(); + ins.type = PTXOperand::f16; + ins.d = f16reg(0); + ins.a = f16reg(1); + ins.b = f16reg(2, true); + setHalf(1, H2); + setHalf(2, H1); + + ins.opcode = PTXInstruction::Add; + cta->eval_Add(cta->getActiveContext(), ins); + if (!checkHalf("add.f16", 0, H3)) return false; + + ins.opcode = PTXInstruction::Sub; + cta->eval_Sub(cta->getActiveContext(), ins); + if (!checkHalf("sub.f16", 0, H1)) return false; + + ins.opcode = PTXInstruction::Mul; + cta->eval_Mul(cta->getActiveContext(), ins); + if (!checkHalf("mul.f16", 0, H2)) return false; + + ins.opcode = PTXInstruction::Min; + cta->eval_Min(cta->getActiveContext(), ins); + if (!checkHalf("min.f16", 0, H1)) return false; + + ins.opcode = PTXInstruction::Max; + cta->eval_Max(cta->getActiveContext(), ins); + if (!checkHalf("max.f16", 0, H2)) return false; + + // abs and neg of -1 + ins = PTXInstruction(); + ins.type = PTXOperand::f16; + ins.d = f16reg(0); + ins.a = f16reg(1); + setHalf(1, HNEG1); + + ins.opcode = PTXInstruction::Abs; + cta->eval_Abs(cta->getActiveContext(), ins); + if (!checkHalf("abs.f16", 0, H1)) return false; + + ins.opcode = PTXInstruction::Neg; + cta->eval_Neg(cta->getActiveContext(), ins); + if (!checkHalf("neg.f16", 0, H1)) return false; + + // fma.rn.f16: 2 * 1 + 1 == 3 + ins = PTXInstruction(); + ins.opcode = PTXInstruction::Fma; + ins.type = PTXOperand::f16; + ins.modifier = PTXInstruction::rn; + ins.d = f16reg(0); + ins.a = f16reg(1); + ins.b = f16reg(2); + ins.c = f16reg(3); + setHalf(1, H2); + setHalf(2, H1); + setHalf(3, H1); + cta->eval_Fma(cta->getActiveContext(), ins); + if (!checkHalf("fma.rn.f16", 0, H3)) return false; + + // set.eq.f16.f16 writes 1.0h or 0.0h, into a register nvrtc declares .b16 + for (int equal = 0; equal < 2; equal++) { + ins = PTXInstruction(); + ins.opcode = PTXInstruction::Set; + ins.type = PTXOperand::f16; + ins.comparisonOperator = PTXInstruction::Eq; + ins.d = f16reg(0, true); + ins.a = f16reg(1, true); + ins.b = f16reg(2, true); + setHalf(1, H2); + setHalf(2, equal ? H2 : H1); + cta->eval_Set(cta->getActiveContext(), ins); + if (!checkHalf("set.eq.f16.f16", 0, equal ? H1 : H0)) return false; + } + + // mov.b16 of a half constant into an .f16 register did nothing at all, + // because eval_Mov_imm switches on d.type and fell through to default + ins = PTXInstruction(); + ins.opcode = PTXInstruction::Mov; + ins.type = PTXOperand::b16; + ins.d = f16reg(0); + ins.a = imm_uint("c", PTXOperand::b16, H4); + setHalf(0, H0); + cta->eval_Mov(cta->getActiveContext(), ins); + if (!checkHalf("mov.b16 immediate", 0, H4)) return false; + + return true; + } + bool test_Abs() { bool result = true; @@ -4445,6 +4603,7 @@ class TestInstructions: public Test { // cvt instruction // arithmetic instructions + result = (result && test_f16()); result = (result && test_Abs()); result = (result && test_Add()); result = (result && test_Sub());