GPU-Unlock: Issue-Rate-Sonde — Rohtexte der Messlaeufe ===================================================== Methode: ein Block pro SM, vier unabhaengige Tensorcore-Instruktionen pro Iteration und Warp (mma.sync.m16n8k16, FP16-Eingaenge, FP32-Akkumulation), Zyklenzahlung ueber das Taktzaehler-Register des SM. Die Sonde ist ein eigenes Programm (C + Hand-PTX, vom Grafiktreamper zur Laufzeit uebersetzt); es braucht kein Rechen-Toolkit der Grafikkartenhersteller. Kernelgroesse 256 und 512 Threads, je beide Karten, 10 Messlaeufe, zwischendurch wechselnde Last auf dem Rechner. Lesart der Zahlen: - "Zyklen pro mma pro SM": wie viele Taktzyklen eine Tensorcore-Instruktion pro Rechenblock braucht. Kleiner ist besser. Die halbe Rate der GeForce-Ausfuehrung dieses Chips liegt bei 8,00, die volle Profi-Rate bei rund 4,00. - "FLOP/SM/Zyklus": Rechenarbeit pro Rechenblock und Taktzyklus, damit unabhaengig von Taktrate und Leistungslimit. Herstellerangabe GeForce-Stock: 510. Volle Profi-Rate desselben Chips (RTX A6000, 165 TF bei 84 Blocks und 1,80 GHz): 1093. - Die FP32-Zeile ist die Eichkurve: sie muss nahe an 256 FLOP/SM/Zyklus landen (128 Rechenwerke pro Block, je zwei Operationen pro Befehl). Bezugswerte, aus der Architekturdokumentation und den Herstellerangaben des Chips: GeForce-Stock 3090: 71,0 TF FP16 mit FP32-Akkumulation = 510 FLOP/SM/Zyklus RTX A6000 (gleicher Chip, volle Rate): 165 TF = 1093 FLOP/SM/Zyklus --- Lauf 1, Blockgroesse 256 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.24 Zyklen/Iteration -> 1014.2 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.2 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.60 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 32.6 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70355..70383) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.22 Zyklen/Iteration -> 1014.3 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.5 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 32.9 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70343..70372) --- Lauf 1, Blockgroesse 512 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.17 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.4 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.34 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 30.8 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149278..149370) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.16 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.6 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.35 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 31.0 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149297..149381) --- Lauf 2, Blockgroesse 256 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.24 Zyklen/Iteration -> 1014.2 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.2 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.60 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 32.6 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70353..70386) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.22 Zyklen/Iteration -> 1014.3 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.5 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 32.9 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70343..70370) --- Lauf 2, Blockgroesse 512 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.18 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.4 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.35 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 30.8 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149267..149382) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.16 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.6 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.34 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 31.0 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149301..149374) --- Lauf 3, Blockgroesse 256 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.24 Zyklen/Iteration -> 1014.2 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.2 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 32.6 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70328..70378) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.22 Zyklen/Iteration -> 1014.3 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.5 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 32.9 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70342..70370) --- Lauf 3, Blockgroesse 512 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.17 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.4 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.35 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 30.8 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149278..149383) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.17 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.6 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.34 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 31.0 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149296..149373) --- Lauf 4, Blockgroesse 256 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.23 Zyklen/Iteration -> 1014.2 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.2 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 32.6 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70331..70380) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.23 Zyklen/Iteration -> 1014.3 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.5 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 32.9 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70342..70374) --- Lauf 4, Blockgroesse 512 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.17 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.4 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.35 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 30.8 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149278..149383) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.16 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.6 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.34 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 31.0 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149299..149373) --- Lauf 5, Blockgroesse 256 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.24 Zyklen/Iteration -> 1014.2 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.2 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 32.6 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70329..70378) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 129.22 Zyklen/Iteration -> 1014.3 FLOP/SM/Zyklus | 4.04 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.5 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 17.59 Zyklen/Iteration -> 232.8 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 32.9 TFLOPS FP32 (91% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 70342..70372) --- Lauf 5, Blockgroesse 512 --- GPU 0 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1710 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.17 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1710 MHz: 142.4 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.35 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1710 MHz: 30.8 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149267..149385) GPU 1 NVIDIA GeForce RTX 3090 SMs 82 max Takt 1725 MHz sm_86 mma.sync.m16n8k16 (Tensor-Core) 258.16 Zyklen/Iteration -> 1015.4 FLOP/SM/Zyklus | 4.03 Zyklen pro mma pro SM hochgerechnet bei 1725 MHz: 143.6 TFLOPS FP16 (FP32-Akkum.) [GeForce-Stock-Spec 3090: 71,0 TF = 510 FLOP/SM/Zyklus] Mehrfach gegen Stock-Spezifikation: 1.99x (volle GA102-Rate, RTX-A6000-Masstab: 1093 FLOP/SM/Zyklus) fma.rn.f32 (FP32-Kerne) 37.34 Zyklen/Iteration -> 219.4 FLOP/SM/Zyklus | FP32-Spitze = 256 FLOP/SM/Zyklus (128 Kerne/SM) hochgerechnet bei 1725 MHz: 31.0 TFLOPS FP32 (86% der Kern-Spitze) (Pruefsumme 853460.4, Zyklen-Streuung 149299..149362) Zusammenfassung --------------- 10 Laeufe, beide Karten, beide Kernelgroessen: 4,03 bis 4,04 Zyklen pro Tensorcore-Instruktion pro Rechenblock, 1014 bis 1015 FLOP pro Rechenblock und Taktzyklus. Hochgerechnet auf den erlaubten Maximaltakt (1710 MHz bzw. 1725 MHz, 82 Rechenblocks) sind das 142,2 bzw. 143,5 TF FP16 mit FP32-Akkumulation — das 1,99-fache der GeForce-Angabe von 71,0 TF und 93 Prozent der vollen Profi-Rate desselben Chips. Streuung ueber alle Laeufe: unter 0,15 Prozent. Eichkurve FP32: 232,8 FLOP/SM/Zyklus = 91 Prozent der Kernspitze von 256.