AlphaGitHub

Hardware / NVIDIA Ampere · GA102

GeForce RTX 3090

Alpha's performance card: Coppelius trains here 1.51× faster than PyTorch with CUDA graphs.

SM86 · compute capability 8.6

What has run here

Qualified

Bob and Coppelius train, resume and predict here. On 27 Sep 2026 Coppelius completed 10,000 steps with every loss finite and resumed in a fresh process, then beat PyTorch in matched 2,000-step runs from the same checkpoint. Performance work happens on this card.

Driver facts

Typed profile

Read at build time from Platform.RunPod.Nvidia.SM86 at 2877a0c1. Systems bind to these values by name; none branches on a card name.

FactValueAlpha definition
Product identityNVIDIA GeForce RTX 3090runPodRTX3090ProductIdentity
Target identityrunpod-nvidia-geforce-rtx3090-sm86runPodRTX3090TargetIdentity
Compute capability8.6runPodRTX3090TargetProfile
Total video memory25,769,803,776 B (24 GiB)runPodRTX3090TotalVideoMemory
QMD architecture0x00000086runPodSM86QMDArchitecture
Compute class0x0000C7C0runPodSM86ComputeClass
SPA version0x00000806runPodSM86SPAVersion

Instruction set · 40 typed instructions

Every GPU instruction Alpha can emit for this device

The typed instruction set is SM86InstructionBody in Accelerator.SM86.Instruction; each name links to the line that defines it, each opcode to the line that encodes it. The SASS column is the mnemonic NVIDIA's tools would print for the same encoding. Instructions this device has that Alpha has not yet typed are not listed: Alpha only claims what its compiler can emit and its runtime has executed.

Alpha nameSASSOpcodeMeaningOperands
SM86MoveConstantMOV (constant bank)0xA02destination = c[bank][offset], a 32-bit word from a constant bank.
SM86SpecialToRegisterS2R0x919destination = a special register (thread or cooperative-thread-array index).
SM86MoveImmediateMOV (immediate)0x802destination = a 32-bit immediate.
SM86IntegerMultiplyAddConstantIMAD (constant bank)0xA24destination = left * c[bank][offset] + addend, 32-bit integer.
SM86IntegerMultiplyAddImmediateIMAD (immediate)0x824destination = left * immediate + addend, 32-bit integer.
SM86IntegerMultiplyAddWideConstantIMAD.WIDE.U320x625destination pair = left * right + c[bank][offset], a 64-bit product plus a 64-bit constant.
SM86IntegerAddThreeImmediateIADD3 (immediate)0x810destination = left + immediate, the three-input adder with one input zero.
SM86IntegerAddThreeRegisterIADD3 (register)0x210destination = left + right, the three-input adder with one input zero.
SM86ShiftRightImmediateSHF.R.U32.HI0x819destination = source logically shifted right by an immediate amount.
SM86LogicThreeInputTruthTableLOP3.LUT0x212destination = truthTable(left, right, zero), any bitwise function of the inputs given as an 8-bit lookup table.
SM86IntegerToFloatI2FP.F32.S320x245destination = the signed 32-bit integer source converted to binary32.
SM86FloatAddFADD0x221destination = left + right, binary32, round to nearest even.
SM86FloatMultiplyFMUL0x220destination = left * right, binary32, round to nearest even.
SM86FloatFusedMultiplyAddFFMA0x223destination = left * right + addend, binary32 with a single rounding.
SM86MultiFunctionUnitApproximationMUFU0x308destination = an approximate transcendental of the source (cosine, sine, exp2, log2, reciprocal, reciprocal square root, square root, tanh) from the multi-function unit.
SM86FloatMinimumOrMaximumFMNMX0x209destination = the minimum or the maximum of left and right, binary32.
SM86FloatNegateFADD (negated source, zero addend)0x221destination = -source, binary32.
SM86FloatPairToPackedHalfPairF2FP.PACK_AB0x23Edestination = two binary32 sources converted and packed into one register of two binary16 halves.
SM86FloatPairToPackedBFloat16PairF2FP.BF16.PACK_AB—destination = two binary32 sources rounded to bfloat16 (to nearest, ties to even) and packed into one register.
SM86HalfToFloatHADD2.F320x230destination = the selected binary16 half of the source widened to binary32.
SM86BFloat16ToFloatPRMT Rd, Ra, 0x1044 (H0) or 0x3244 (H1), RZ—destination = the selected bfloat16 half of the source widened to binary32 (its bits moved to the high half, the low half zero: exact).
SM86TensorCoreHalfMatrixMultiplyAccumulate16x8x16Float32HMMA.16816.F320x23Cdestination = fragmentA (16x16, binary16) * fragmentB (16x8, binary16) + accumulator (16x8, binary32) on the tensor cores, one warp-wide 16x8x16 tile.
SM86TensorCoreBFloat16MatrixMultiplyAccumulate16x8x16Float32HMMA.16816.F32.BF16—the same tile product with bfloat16 fragments.
SM86LoadGlobalLDG.E0x981destination = the 32-bit word at address + offset in global memory.
SM86LoadGlobalWideLDG.E.1280x981destination quad = the 128 bits at address + offset in global memory.
SM86WarpShuffleSHFL0xF89destination = the source register of another lane of the warp, selected by index, up, down or butterfly mode.
SM86LoadSharedLDS0x984destination = the 32-bit word at address + offset in shared memory.
SM86LoadSharedMatrixLDSM0x83Bdestination = one, two or four 8x8 binary16 matrices loaded from shared memory into tensor-core fragment layout, optionally transposed.
SM86StoreSharedSTS0x388the 32-bit value stored at address + offset in shared memory.
SM86LoadGlobalToSharedLDGSTS.E.BYPASS.128 (cp.async.cg, 16 bytes)0xFAEthe 128 bits at global source + sourceOffset (a register pair; the offset a signed 12-bit immediate) copied to shared memory at address + offset, asynchronously: the copy joins the thread's open group (SM86CommitAsyncGroup closes it), and nothing may read those shared bytes until SM86WaitAsyncGroups has retired the group (and, for other threads, a BAR.SYNC after). Its registers are read after issue: a read barrier guards them.
SM86CommitAsyncGroupLDGDEPBAR (cp.async.commit_group)0x9AFthe thread's open group of asynchronous copies closed; the groups in flight count on scoreboard SB0, which the control's write barrier names.
SM86WaitAsyncGroupsDEPBAR.LE SB0, count (cp.async.wait_group count)0x91Await until at most `count` (0..63) groups of asynchronous copies are in flight.
SM86BarrierSynchronizeBAR.SYNC0xB1Devery thread of the cooperative thread array waits until all have arrived.
SM86StoreGlobalSTG.E0x986the 32-bit value stored at address + offset in global memory.
SM86StoreGlobalWideSTG.E.1280x986the 128-bit value quad stored at address + offset in global memory.
SM86StoreGlobal64STG.E.640x986the 64-bit value pair stored at address + offset in global memory.
SM86ReduceGlobalAddFloat32RED.E.ADD.F320x98Ethe binary32 value atomically added to the word at address + offset in global memory, no result returned.
SM86PredicateGreaterThanImmediateISETP.GT (immediate)0x80Cdestination predicate = source > immediate, compared as unsigned 32-bit words in this encoding: the RTX 3070 has 0x80000000 > 0 (Checked.IntegerCompareProbe, Proof.CheckedIntegerCompareSilicon).
SM86BranchBRA0x947continue at the instruction offset bytes away, relative to the next instruction.
SM86ExitEXIT0x94Dthis thread finishes the program.

Operand vocabularies

The named values an instruction's operands range over: special registers, execution guards, shuffle modes, half selectors, shared-matrix shapes, multi-function operations, extremum modes, and the scoreboard barriers and yield modes carried by every instruction's SM86Control word (Stall, Yield, WriteBarrier, ReadBarrier, WaitMask, ReuseMask).

SM86SpecialRegister

Alpha nameSASSMeaning
SM86CooperativeThreadArrayIdXSR_CTAID.Xthe cooperative thread array (block) index on x.
SM86CooperativeThreadArrayIdYSR_CTAID.Ythe cooperative thread array (block) index on y.
SM86CooperativeThreadArrayIdZSR_CTAID.Zthe cooperative thread array (block) index on z.
SM86ThreadIdXSR_TID.Xthe thread index on x within its cooperative thread array.
SM86ClockLowSR_CLOCKLOthe low word of the SM's cycle counter.
SM86GlobalTimerLowSR_GLOBALTIMERLOthe low word of the device's global timer, in nanoseconds -- the clock the semaphore releases stamp.
SM86GlobalTimerHighSR_GLOBALTIMERHIits high word.

SM86InstructionGuard

Alpha nameSASSMeaning
SM86InstructionAlways—no guard: the instruction always executes.
SM86InstructionWhen@Pthe instruction executes only in threads where the predicate is true.
SM86InstructionWhenNot@!Pthe instruction executes only in threads where the predicate is false.

SM86ShuffleMode

Alpha nameSASSMeaning
SM86ShuffleIndexSHFL.IDXread the lane named by the lane operand.
SM86ShuffleUpSHFL.UPread the lane that is lane places below this one.
SM86ShuffleDownSHFL.DOWNread the lane that is lane places above this one.
SM86ShuffleButterflySHFL.BFLYread the lane whose index is this lane's index xor the lane operand.

SM86HalfSelector

Alpha nameSASSMeaning
SM86LowHalf—the low binary16 half (H0) of a 32-bit register.
SM86HighHalf—the high binary16 half (H1) of a 32-bit register.

SM86HalfFormat

Alpha nameSASSMeaning
SM86Binary16——
SM86BFloat16——

SM86SharedMatrixCount

Alpha nameSASSMeaning
SM86SharedMatrix1LDSM.x1one 8x8 matrix.
SM86SharedMatrix2LDSM.x2two 8x8 matrices.
SM86SharedMatrix4LDSM.x4four 8x8 matrices.

SM86SharedMatrixTranspose

Alpha nameSASSMeaning
SM86SharedMatrixNotTransposedLDSM (no .T)rows as stored.
SM86SharedMatrixTransposedLDSM.Teach matrix transposed while loading.

SM86MultiFunction

Alpha nameSASSMeaning
SM86CosineMUFU.COScosine of the source, radians.
SM86SineMUFU.SINsine of the source, radians.
SM86ExponentialBase2MUFU.EX2two raised to the source.
SM86LogarithmBase2MUFU.LG2the base-2 logarithm of the source.
SM86ReciprocalMUFU.RCPone divided by the source.
SM86ReciprocalSquareRootMUFU.RSQone divided by the square root of the source.
SM86SquareRootMUFU.SQRTthe square root of the source.
SM86HyperbolicTangentMUFU.TANHthe hyperbolic tangent of the source.

SM86FloatExtremum

Alpha nameSASSMeaning
SM86FloatMinimumFMNMX with PTthe smaller of the two operands.
SM86FloatMaximumFMNMX with !PTthe larger of the two operands.

SM86YieldMode

Alpha nameSASSMeaning
SM86Continue——
SM86Yield——

Profile coverage

What has been measured, and what has not

A fully profiled device has three tiers. A tier that is not measured yet lists the fields it will fill, so the gap is explicit.

  1. 1

    Driver facts

    Live

    What the resource manager reports and the typed profile binds: product identity, compute capability, video memory, the QMD architecture, compute class and SPA version the command stream is built for.

    • product identity
    • compute capability
    • total video memory
    • QMD architecture
    • compute class
    • SPA version
    • UVM / RM / QMD / pushbuffer / GPFIFO admission bits
  2. 2

    Micro-benchmarks

    Not measured yet

    Measured instruction and memory behaviour, so the compiler can reason about cost from numbers read on this card rather than assumptions.

    • issue rate and latency per typed instruction (FFMA, IMAD, MUFU, LDG, STG, LDSM, HMMA…)
    • shared-memory bandwidth and bank behaviour
    • global-memory bandwidth by access pattern and width
    • tensor-core throughput at 16×8×16
    • barrier and warp-shuffle cost
    • submission, doorbell and semaphore round-trip cost
    • sustained clocks under the training workload
  3. 3

    Model-shaped calibration

    First reading

    The cost of the exact executables Alpha emits for each system on this device. End-to-end throughput and per-launch device times from the GPU's own timestamps exist (below); occupancy and register pressure have not been measured.

    • per-kernel time for every launch in each system's schedule
    • occupancy and register / shared-memory pressure per kernel
    • end-to-end step time and its split between host protocol, launches and waits
    • memory footprint against the profile's video memory

Measured on this card · 2026-09-27

Measured readings

From the run record. These measure Alpha's current executables on this card, not the card's limits.

ReadingValue
Coppelius training, matched 2,000-step runs (wall, checkpoints included)Alpha 35,683 tokens/s · PyTorch CUDA graphs 23,675 · PyTorch eager 15,180
Mean loss over those runsAlpha 6.8840 · PyTorch 6.8818 (largest 1,000-step block difference 0.0038 nats)
10,000-step run, ten checkpoints included36,310 tokens/s wall; every loss finite; fresh-process resume completed
One profiled training step (658 launches)27.4 ms on the device's own clock
243 matrix products in one stepAlpha 11.05 ms · PyTorch 8.54 ms
AdamW update in one stepAlpha 2.05 ms · PyTorch 20.90 ms
Driver580.126.20

Sources

Pairings and evidence

Typed profileBob pairingCoppelius pairingAlpha vs PyTorch recordOriginal qualification reportSystems that target this device