Gibt es eine Möglichkeit, die Anzahl der Taktzyklen zu finden, die für verschiedene Operationen wie Division, Subtraktion und Addition in GPU mit CUDA benötigt werden?Anzahl der Taktzyklen pro Operation in GPU
Dank
Sie vielleicht einen Blick auf die Uhr Gerätefunktion haben wollen: clock. Es kann nicht eine Grundwahrheitsindikator der Zyklen pro Anweisung sein, die Sie erwähnen, wie Latenz und Throuhput ins Spiel kommen, aber ist definitiv ein Werkzeug, das helfen wird.
Hier ein Codebeispiel, wie es für Durchsatzschätzung verwendet werden könnte:
__global__ void timing_1(int N, float4* input, float4* output)
{
float4 a,b,c,d ;
a = input[0]; b = input[1]; c = input[2]; d = input[3];
long long int start = clock64();
for (int k = 0 ; k < N ; ++k)
{
a.x = 1.0/a.x ; a.y = 1.0/a.y ; a.z = 1.0/a.z ; a.w = 1.0/a.w ;
b.x = 1.0/b.x ; b.y = 1.0/b.y ; b.z = 1.0/b.z ; b.w = 1.0/b.w ;
c.x = 1.0/c.x ; c.y = 1.0/c.y ; c.z = 1.0/c.z ; c.w = 1.0/c.w ;
d.x = 1.0/d.x ; d.y = 1.0/d.y ; d.z = 1.0/d.z ; d.w = 1.0/d.w ;
}
long long int stop = clock64();
// make use of data so that compiler does not optimize it out
a.x += b.x + c.x + d.x ;
a.y += b.y + c.y + d.y ;
a.z += b.z + c.z + d.z ;
a.w += b.w + c.w + d.w ;
output[threadIdx.x + blockDim.x * blockIdx.x] = a ;
if (threadIdx.x == 0)
::printf ("timing_1 - Block [%d] - cycles count = %lf - cycles per div = %lf\n", blockIdx.x, ((double)(stop - start)), ((double)(stop-start))/(16.0*(double)N)) ;
}
Für Latenz, möchten Sie die Abhängigkeit zwischen Berechnungen haben:
__global__ void timing_2(int N, float4* input, float4* output)
{
float4 a ;
a = input[0];
long long int start = clock64();
for (int k = 0 ; k < N ; ++k)
{
a.y = 1.0/a.x ; a.z = 1.0/a.y ; a.w = 1.0/a.z ; a.x = 1.0/a.w ;
}
long long int stop = clock64();
output[threadIdx.x + blockDim.x * blockIdx.x] = a ;
if (threadIdx.x == 0)
::printf ("timing_2 - Block [%d] - cycles count = %lf - cycles per div = %lf\n", blockIdx.x, ((double)(stop - start)), ((double)(stop-start))/(4.0*(double)N)) ;
}
Sie wollen dies mit einem laufen kleine Anzahl von Threads und Blöcken pro SM, um Überlappungen von Berechnungen zu vermeiden, die Ihren Wanduhr-Timer mit einzelnen Berechnungen inkonsistent machen würden.
Für 32 Threads und 5 Blöcke auf einer GTX 850m, erhalte ich einen Durchsatz von 128 Zyklen pro Division und eine Latenz von 142 Zyklen in einfacher Genauigkeit mit normaler Mathematik - verwandelt sich in einen Funktionsaufruf (nvcc 7.5, sm_50). 
Bei der Verwendung von Fast-Math, bekomme ich einen Durchsatz von 2,5 Zyklen und Latenz von 3 Zyklen. 
Bitte erarbeiten Sie eine tatsächliche Antwort, nicht nur externe Links – abarisone
Ich glaube, diese Antwort ist Borderline irreführend, da es sehr schwierig wäre, Instruction-Level-Timings mit dieser Funktion zu messen. Zumindest sollte diese Antwort mit einem Micro-Benchmarking-Code oder einer Art Erklärung dafür, wie man es erstellt, kommen. –
in der Tat, mit '--use_fast_math' auf einem Titan X (Maxwell, Rechenleistung 5.2) bekomme ich einen Durchsatz von 0,69 und eine Latenz von 0,66 Zyklen pro Division (wenn ich' --use_fast_math' weglasse, ist der Durchsatz viel kleiner und die Latenz viel höher). In der schnellen mathematischen Version ist die erzeugte ptx-Anordnung voll von "rcp.approx" -Befehlen, d. H. Einer dedizierten Anweisung zum Berechnen von reziproken Werten. –
Diese Frage könnte entweder * Latenz * oder * Durchsatz * im Blick haben. Durchsätze (in ops/clock) sind in [dem Programmierleitfaden] (http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#arithmetic-instructions) für eine Anzahl von Operationen dokumentiert. Latenzen sind meistens undokumentierte AFAIK, jedoch gibt es einige Mikrobenchmarking-Übungen (wie [diese] (http://www.eecg.toronto.edu/~myrto/gparch-ispass2010.pdf)), die durchgeführt wurden, um Latenzen zu bestimmen . Latenzen (und Durchsätze) variieren je nach GPU-Architektur. –
GPUs bieten keine Hardware-Anweisungen für die Division. Sowohl die Ganzzahldivision als auch die Gleitkommadivision sind Softwareroutinen, die Inline-Routinen, teilweise Inline-Routinen oder aufgerufene Subroutinen sein können. Da es sich um Software handelt, kann * sie sich mit der CUDA-Version ändern und * ändert sich * mit der GPU-Architektur. Für Gleitkommazahlen mit einfacher Genauigkeit stehen mehrere Varianten zur Verfügung, die sich in Leistung und Genauigkeit unterscheiden. Sowohl die Gleitkomma-Division mit einfacher Genauigkeit als auch mit doppelter Genauigkeit haben Varianten mit unterschiedlichen Rundungsmodi und unterschiedlicher Leistung. Die Integer-Teilungsgeschwindigkeit unterscheidet sich durch Bitbreite und Vorzeichen/Vorzeichen. – njuffa
@njuffra, was ist Ihre Ansicht auf __fdividef wie hier dargestellt http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#intrinsic-functions, ist dies intrinsische Exposition Special Function Unit aus Umfang ? –