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/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/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/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()); diff --git a/ocelot/src/ir/PTXInstruction.cpp b/ocelot/src/ir/PTXInstruction.cpp index 099a0493..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 ); } @@ -667,12 +666,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 ); @@ -921,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 ); } @@ -945,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 ) ) { @@ -1122,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 ); } @@ -1152,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 ); } @@ -1178,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 ); } @@ -1222,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 ); } @@ -1263,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 ); } @@ -1575,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 ) { @@ -1614,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 ); } @@ -1862,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 ); } 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; 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(); +}