ComputerBase Menü
Aktuelles

3DCenter berichtet über angeblichen VRAM Fehler der 970

Status
Für weitere Antworten geschlossen.
A

AMINDIA

Banned
Registriert
Juli 2012
Beiträge
1.382
Wäre schön, wenn CB mal dieser Sache nachgehen würde.
 
@Nai

An sich braucht man einen Benchmark, dem die reine Geschwindigkeit der Karte egal ist, der aber den VRam innerhalb einer bestimmten Zeit immer weiter befüllt. Wenn Du etwas in dieser Art schreiben kannst wäre das echt super....:)
 
Mir schwebte eher soetwas vor, was sich nacheinander Speicherblöcke auf der GPU alloziert bis kein Speicher mehr frei ist, und dann innerhalb von jedem dieser Blöcke die DRAM- und Cache-Bandbreite misst. Oder alternative Vorschläge ?

Der Vorteil an CUDA ist halt, dass es im Vergleich zu OpenCL(?) DirectX oder OpenGL kein automatisches Speichermanagement besitzt. Deshalb kann die GPU den CUDA-Speicher nicht in den CPU-DRAM auslagern. Dadurch könnte man ausschließen ob es am Swapping liegt, sondern zeigen, dass es wirklich etwas mit dem Caching oder der DRAM-Anbindung zu tun hätte.

Edit:
So habs eben mal aus "langweile" programmiert. Irgendjemand vorschläge wo ich heutzutäge möglichst einfach eine Exe hochladen kann ? :)

Fazit des Programmierens: Beobachte das selbe interessanterweise auch auf meiner Titan. Die letzten ~250 MB sind deutlich langsamer als der Rest. Sowohl die Cache-Bandbreite als auch die DRAM-Bandbreite
Ergänzung ()

Hier die Exe:
Rec.exe 16 KB
https://mega.co.nz/#!5pU0TZzI!bYaoUsJN9NvEWcO-2_ocg6cZ5EnrRoLzGqI2HG9mZG0

Hoffe Link ist erlaubt . . . . und hoffe die exe geht so . . . Für weitere Benchmarkwünsche bin ich auch gerne offen.
 
Zuletzt bearbeitet:
  • Gefällt mir
Reaktionen: Vitche, w0mbat, Jan und eine weitere Person
Puh, da schlagen alle Alarmsysteme an, wenn man das testen möchte ^^. Bereits Chrome lässt mich die Datei nicht runterladen.

Deaktiviere ich den Phishing-Malware Schutz, schlägt sofort der Kaspersky Alarm...beende ich den, lässt es der Defender nicht zu ^^.
 
Tut mir leid, das ist eigentlich eine *ganz normale* Visual Studio exe. Mich würde es wirklich extrem überraschen, wenn etwas an der exe faul ist. Hier der Quellcode, falls sie jemand selbst kompilieren möchte:
Code: 
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "helper_math.h"
#include <stdio.h>
#include <iostream>
#define CacheCount 3
__global__ void BenchMarkDRAMKernel(float4* In)
{
	int ThreadID = blockDim.x *blockIdx.x + threadIdx.x ;

	float4 Temp = make_float4(1);

	Temp += In[ThreadID];
	

	if (length(Temp) == -12354)
		In[0] = Temp;

}




__global__ void BenchMarkCacheKernel(float4* In, int Zero)
{
	int ThreadID = blockDim.x *blockIdx.x + threadIdx.x;

	float4 Temp = make_float4(1);

#pragma unroll
	for (int i = 0; i < CacheCount; i++)
	{
		Temp += In[ThreadID + i*Zero];
	}

	if (length(Temp) == -12354)
		In[0] = Temp;

}



int main()
{
	static const int PointerCount = 5000;

	int Float4Count = 8 * 1024 * 1024;
	int ChunkSize = Float4Count*sizeof(float4);
	float4* Pointers[PointerCount];
	int UsedPointers = 0;
	printf("Nai's Benchmark \n");
	printf("Allocating Memory . . . Chunk Size = %i Byte  \n", ChunkSize);
	system("pause");
	while (true)
	{

		int Error = cudaMalloc(&Pointers[UsedPointers], ChunkSize);

		if (Error == cudaErrorMemoryAllocation)
			break;

		cudaMemset(Pointers[UsedPointers], 0, ChunkSize);
		UsedPointers++;
	}
	printf("Allocated %i Chunks \n", UsedPointers);
	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	int BlockSize = 128;
	int BlockCount = Float4Count / BlockSize;

	int BenchmarkCount = 10;

	printf("Benchmarking DRAM \n");
	system("pause");
	for (int i = 0; i < UsedPointers; i++)
	{
		cudaEventRecord(start);
		for (int j = 0; j < BenchmarkCount; j++)
			BenchMarkDRAMKernel <<<BlockCount, BlockSize >>>(Pointers[i]);

		cudaEventRecord(stop);
		cudaEventSynchronize(stop);

		float milliseconds = 0;
		cudaEventElapsedTime(&milliseconds, start, stop);

		float Bandwidth = ((float)(BenchmarkCount * ChunkSize)) / milliseconds/ 1000.f/1000.f;
		printf("DRAM-Bandwidth of %i. Chunk: %f GByte/s \n", i, Bandwidth);
	}


	system("pause");
	printf("Benchmarking L2-Cache \n");
	system("pause");


	for (int i = 0; i < UsedPointers; i++)
	{
		cudaEventRecord(start);
		for (int j = 0; j < BenchmarkCount; j++)
			BenchMarkCacheKernel << <BlockCount, BlockSize >> >(Pointers[i], 0);

		cudaEventRecord(stop);
		cudaEventSynchronize(stop);

		float milliseconds = 0;
		cudaEventElapsedTime(&milliseconds, start, stop);

		float Bandwidth = (((float)CacheCount* (float)BenchmarkCount * (float)ChunkSize)) / milliseconds / 1000.f / 1000.f;
		printf("L2-Cache-Bandwidth of %i. Chunk: %f GByte/s  \n", i, Bandwidth);
	}


	system("pause");

	cudaDeviceSynchronize();
	cudaDeviceReset();
    return 0;
}
Ergänzung ()

Fast vergessen zu posten: Die Ergebnisse von meiner Titan:
Nai's Benchmark
Allocating Memory . . .
Chunk Size = 128 MiByte
Press any key to continue . . .
Allocated 46 Chunks
Allocated 5888 MiByte
Benchmarking DRAM
Press any key to continue . . .
DRAM-Bandwidth of Chunk no. 0 (0 MiByte to 128 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 1 (128 MiByte to 256 MiByte): 249 GByte/s
DRAM-Bandwidth of Chunk no. 2 (256 MiByte to 384 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 3 (384 MiByte to 512 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 4 (512 MiByte to 640 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 5 (640 MiByte to 768 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 6 (768 MiByte to 896 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 7 (896 MiByte to 1024 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 8 (1024 MiByte to 1152 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 9 (1152 MiByte to 1280 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 10 (1280 MiByte to 1408 MiByte): 255 GByte/s
DRAM-Bandwidth of Chunk no. 11 (1408 MiByte to 1536 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 12 (1536 MiByte to 1664 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 13 (1664 MiByte to 1792 MiByte): 249 GByte/s
DRAM-Bandwidth of Chunk no. 14 (1792 MiByte to 1920 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 15 (1920 MiByte to 2048 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 16 (2048 MiByte to 2176 MiByte): 249 GByte/s
DRAM-Bandwidth of Chunk no. 17 (2176 MiByte to 2304 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 18 (2304 MiByte to 2432 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 19 (2432 MiByte to 2560 MiByte): 246 GByte/s
DRAM-Bandwidth of Chunk no. 20 (2560 MiByte to 2688 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 21 (2688 MiByte to 2816 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 22 (2816 MiByte to 2944 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 23 (2944 MiByte to 3072 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 24 (3072 MiByte to 3200 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 25 (3200 MiByte to 3328 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 26 (3328 MiByte to 3456 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 27 (3456 MiByte to 3584 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 28 (3584 MiByte to 3712 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 29 (3712 MiByte to 3840 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 30 (3840 MiByte to 3968 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 31 (3968 MiByte to 4096 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 32 (4096 MiByte to 4224 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 33 (4224 MiByte to 4352 MiByte): 256 GByte/s
DRAM-Bandwidth of Chunk no. 34 (4352 MiByte to 4480 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 35 (4480 MiByte to 4608 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 36 (4608 MiByte to 4736 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 37 (4736 MiByte to 4864 MiByte): 255 GByte/s
DRAM-Bandwidth of Chunk no. 38 (4864 MiByte to 4992 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 39 (4992 MiByte to 5120 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 40 (5120 MiByte to 5248 MiByte): 250 GByte/s
DRAM-Bandwidth of Chunk no. 41 (5248 MiByte to 5376 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 42 (5376 MiByte to 5504 MiByte): 257 GByte/s
DRAM-Bandwidth of Chunk no. 43 (5504 MiByte to 5632 MiByte): 255 GByte/s
DRAM-Bandwidth of Chunk no. 44 (5632 MiByte to 5760 MiByte): 29 GByte/s
DRAM-Bandwidth of Chunk no. 45 (5760 MiByte to 5888 MiByte): 16 GByte/s
Press any key to continue . . .
Benchmarking L2-Cache
Press any key to continue . . .
L2-Cache-Bandwidth of Chunk no. 0 (0 MiByte to 128 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 1 (128 MiByte to 256 MiByte): 370 GByte/s
L2-Cache-Bandwidth of Chunk no. 2 (256 MiByte to 384 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 3 (384 MiByte to 512 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 4 (512 MiByte to 640 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 5 (640 MiByte to 768 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 6 (768 MiByte to 896 MiByte): 374 GByte/s
L2-Cache-Bandwidth of Chunk no. 7 (896 MiByte to 1024 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 8 (1024 MiByte to 1152 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 9 (1152 MiByte to 1280 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 10 (1280 MiByte to 1408 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 11 (1408 MiByte to 1536 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 12 (1536 MiByte to 1664 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 13 (1664 MiByte to 1792 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 14 (1792 MiByte to 1920 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 15 (1920 MiByte to 2048 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 16 (2048 MiByte to 2176 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 17 (2176 MiByte to 2304 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 18 (2304 MiByte to 2432 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 19 (2432 MiByte to 2560 MiByte): 370 GByte/s
L2-Cache-Bandwidth of Chunk no. 20 (2560 MiByte to 2688 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 21 (2688 MiByte to 2816 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 22 (2816 MiByte to 2944 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 23 (2944 MiByte to 3072 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 24 (3072 MiByte to 3200 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 25 (3200 MiByte to 3328 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 26 (3328 MiByte to 3456 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 27 (3456 MiByte to 3584 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 28 (3584 MiByte to 3712 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 29 (3712 MiByte to 3840 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 30 (3840 MiByte to 3968 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 31 (3968 MiByte to 4096 MiByte): 370 GByte/s
L2-Cache-Bandwidth of Chunk no. 32 (4096 MiByte to 4224 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 33 (4224 MiByte to 4352 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 34 (4352 MiByte to 4480 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 35 (4480 MiByte to 4608 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 36 (4608 MiByte to 4736 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 37 (4736 MiByte to 4864 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 38 (4864 MiByte to 4992 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 39 (4992 MiByte to 5120 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 40 (5120 MiByte to 5248 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 41 (5248 MiByte to 5376 MiByte): 377 GByte/s
L2-Cache-Bandwidth of Chunk no. 42 (5376 MiByte to 5504 MiByte): 373 GByte/s
L2-Cache-Bandwidth of Chunk no. 43 (5504 MiByte to 5632 MiByte): 370 GByte/s
L2-Cache-Bandwidth of Chunk no. 44 (5632 MiByte to 5760 MiByte): 16 GByte/s
L2-Cache-Bandwidth of Chunk no. 45 (5760 MiByte to 5888 MiByte): 30 GByte/s
Press any key to continue . . .
 
Zuletzt bearbeitet:
Angefangen hat es wohl hier:

http://forums.guru3d.com/showthread.php?t=396064

Keiner weiß bis jetzt was da los ist und ob da überhaupt was los ist. Jedenfalls berichten viele Leute von massiven Framedrops sobald die 3,5GB Speicher überschritten werden, bzw. sie werden gar nicht überschritten. Sehr merkwürdig das ganze und keiner hat bisher eine Erklärung dafür finden können.

Mal sehen wie sich das entwickelt.
 
Bei mir kommt immer folgender Fehler...
 

Anhänge

  • problem_cudart.JPG
    problem_cudart.JPG
    25,2 KB · Aufrufe: 1.960
Zuletzt bearbeitet:
Was ist denn die Kernaussage von dem was ich hier sehe?


Ich deute es so, dass es bereits ab 3GB einen erheblichen Einbruch in der Speicherbandbreite gibt. Ich habe auch eben mal den Speicher ordentlich übertaktet (+400 Mhz) und die Werte haben sich nur minimal verändert. Leicht verbessert, aber teilweise auch verschlechtert.
 

Anhänge

  • nais_bench_1.JPG
    nais_bench_1.JPG
    281 KB · Aufrufe: 2.188
Zuletzt bearbeitet:
Die Kernaussage ist folgende:
Diejenigen 3 Gigabyte, die dir CUDA zuerst alloziert (= reserviert), besitzen eine DRAM-Bandbreite von ~ 155 Gigabyte/s
Dasjenige Gigabyte, das dir CUDA zuletzt alloziert, besitzt eine DRAM-Bandbreite von ~ 22 Gigabyte/s
Mit der Cache-Bandbreite verhält es sich analog. Wieso das so ist, kann ich leider keine Aussage treffen.

Auf meiner Titan sind es in etwa nur die letzten ~250 MByte die ähnlich langsam sind
 
Also bedeutet das, dass 0-3 Gigabyte mit voller Geschwindigkeit angesprochen werden können, und es ab dem 3. - 4. Gigabyte einen massiven Einbruch gibt.

Dieser beträgt bei deiner Titan 256MB und bei der GTX970 satte 1024MB, ist das so korrekt?
 
Also bedeutet das, dass 0-3 Gigabyte mit voller Geschwindigkeit angesprochen werden können, und es ab dem 3. - 4. Gigabyte einen massiven Einbruch gibt.
Zum Vergrößern anklicken....

Ja das ist korrekt.

Dieser beträgt bei deiner Titan 256MB und bei der GTX970 satte 1024MB, ist das so korrekt?
Zum Vergrößern anklicken....
Hier muss ich mich etwas selbst berichtigen; ich habe noch etwas herumgetüftelt. Die Einbrüche scheinen durch das Swapping zu passieren, also dadurch dass die GPU ihren GPU-DRAM in den CPU-DRAM auslagert. Denn es haben in so viel Megabyte die geringere Bandbreite, wie du vor dem Benchmark-Programm bereits belegt hattest. Interessanterweise scheint er nur immer den als letztes allozierten Speicher zu swappen.

Auch scheint er den geswapten Speicher nicht wieder in den GPU-DRAM hochzuladen, sondern weiter immer direkt auf den CPU-DRAM zuzugreifen. Deshalb ist auch die *geschätzte* Cache-Bandbreite so schlecht.
 
Irgendwas stimmt mit deinem Code noch nicht ganz. Ich komme auf z.B. auf 1048576 GByte/s. Das wäre eine sagenhafte Datendurchsatzrate.



Wenn ich deinen Code richtig lese allozierst du Chunk * Block und kopierst dann Nullen rüber und lässt nen counter laufen. Wann stoppt dieser denn? Wo finde ich denn cudaEventRecord(stop)? Ansonsten sehr gut dass du dich der Sache angenommen hast. Komm doch zu uns ins 3DC :)

Meine Karten sind zwei GTX 770 mit je 4GB.

ps: Bitte nehme die Pausen bis auf die letzte heraus. Oder eine Version ganz ohne, damit man ne Schleife basteln kann. Hier auf den PC kommt kein Visual Basic rauf und auf dem Notebook hab ichs noch nicht installiert :D

Edit: Ach ja und kennst du eine Möglichkeit die größe des L2$ heraus zu bekommen? Gibt ja keine MSR wie bei einer CPU die das mitteilen.
 
Zuletzt bearbeitet:
Vergleicht man die Werte mit theoretischen Peaks kommt das relativ gut hin. Nur auf GK104 definitiv nicht :D Exklusiver Fenstermodus wäre doch ne Alternative? Oder halt für bash compilen.
 
Im Fenstermodus kann das wohl nicht präzise sein.
Zum Vergrößern anklicken....
Bis zum Swapping ist es doch relativ präzise, wie man an den Messergebnissen sieht. Und auch die Messungen beim Swapping sind "präzise", wenn man sie richtig interpretiert. Selbst wenn das Swapping in CUDA verglichen mit OpenGL und DirektX vermutlich etwas anders abläuft, weshalb das für euch weniger interessant ist, ist es für mich im Speziellen sehr interessant wie CUDA das macht. Denn ich kann daraus folgern, dass er die Daten einfach assoziativ swapt; und zudem die geswapten CUDA-Speicherbereiche wahrscheinlich nicht mehr sofort auf die GPU hochlädt. Aus den Gründen besitzen immer die selben geswappten Speicherbereiche eine extrem niedrige Bandbreite. Weitere Benchmarks mit ähnlichen Problemfällen in OpenGL oder DirectX wären interessent, um zu sehen wie NVIDIA da das Swapping macht.

Aber auch erkennt man, ob die Speicherbandbreite und die Cache-Bandbreite, sofern kein Swapping eintritt, in allen alllozierten Speicherbereichen in etwa konstant ist oder einbricht. Denn das Einbrechen war ja eine Vermutung für die schlechte Performance der Geforce 970 GTX. Ebenso wird die Cache-Bandbreite in allen Bereichen gemessen, deren Einbruch ja auch eine Vermutung darstellte.

Wenn ich deinen Code richtig lese allozierst du Chunk * Block und kopierst dann Nullen rüber und lässt nen counter laufen. Wann stoppt dieser denn? Wo finde ich denn cudaEventRecord(stop)? Ansonsten sehr gut dass du dich der Sache angenommen hast. Komm doch zu uns ins 3DC
Zum Vergrößern anklicken....

Der Code ist ein relativ großer "Hack". Ich habe die Nullen da nur präventiv reinkopiert, damit das ganze nicht eventuell Sonderbehandlung wegen Denormalized-Floats benötigt. In den Schleifen rufe ich für jeden Chunk mehrmals ein Programm beziehungsweise ein Kernel auf der GPU auf, was einfach mit einem Regelmäßigen Zugriffsmuster jede Cache-Line einmal lädt, wodurch ich die Speicherbandbreite ermittel. Das Start- und Stop-Event dienen dazu die Laufzeiten des Kernels auf der GPU präzise und ohne Overhead zu bestimmen. Bei dem Benchmark für die Cache-Bandbreite lade ich jede Cache-Line mehrmals in folge rein, damit ich die Cache-Bandbreite bestimme.

Die kaputten Laufzeiten kommen bei dir daher, dass ich nur CUDA-Code für Titan und aufwärts in der Exe beigelgt habe. Das und deine Anderen Wünsche und ein paar Kleinigkeiten habe ich in ner neuen Version eben mal gefixt:
Rec.exe 82 KB
https://mega.co.nz/#!JscwXTDS!SZd23Z3wswoqknKTIjNnsHU7lpI_-5YpOw0GdJuzlgI

Code: 
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "helper_math.h"
#include <stdio.h>
#include <iostream>
#define CacheCount 5
__global__ void BenchMarkDRAMKernel(float4* In, int Float4Count)
{
	int ThreadID = (blockDim.x *blockIdx.x + threadIdx.x) % Float4Count;

	float4 Temp = make_float4(1);

	Temp += In[ThreadID];
	

	if (length(Temp) == -12354)
		In[0] = Temp;

}


__global__ void BenchMarkCacheKernel(float4* In, int Zero,int Float4Count)
{
	int ThreadID = (blockDim.x *blockIdx.x + threadIdx.x) % Float4Count;

	float4 Temp = make_float4(1);

#pragma unroll
	for (int i = 0; i < CacheCount; i++)
	{
		Temp += In[ThreadID + i*Zero];
	}

	if (length(Temp) == -12354)
		In[0] = Temp;

}



int main()
{
	static const int PointerCount = 5000;

	int Float4Count = 8 * 1024 * 1024;
	int ChunkSize = Float4Count*sizeof(float4);
	int ChunkSizeMB = (ChunkSize / 1024) / 1024;
	float4* Pointers[PointerCount];
	int UsedPointers = 0;
	printf("Nai's Benchmark \n");
	printf("Allocating Memory . . . \nChunk Size: %i MiByte  \n", ChunkSizeMB);

	while (true)
	{

		int Error = cudaMalloc(&Pointers[UsedPointers], ChunkSize);

		if (Error == cudaErrorMemoryAllocation)
			break;

		cudaMemset(Pointers[UsedPointers], 0, ChunkSize);
		UsedPointers++;
	}


	printf("Allocated %i Chunks \n", UsedPointers);

	printf("Allocated %i MiByte \n", ChunkSizeMB*UsedPointers);

	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	int BlockSize = 128;


	int BenchmarkCount = 30;
	int BlockCount = BenchmarkCount * Float4Count / BlockSize;

	printf("Benchmarking DRAM \n");

	for (int i = 0; i < UsedPointers; i++)
	{
		cudaEventRecord(start);

		BenchMarkDRAMKernel << <BlockCount, BlockSize >> >(Pointers[i], Float4Count);

		cudaEventRecord(stop);
		cudaEventSynchronize(stop);

		float milliseconds = 0;
		cudaEventElapsedTime(&milliseconds, start, stop);

		float Bandwidth = ((float)(BenchmarkCount)* (float)(ChunkSize)) / milliseconds / 1000.f / 1000.f;
		printf("DRAM-Bandwidth of Chunk no. %i (%i MiByte to %i MiByte):%5.2f GByte/s \n", i, ChunkSizeMB*i, ChunkSizeMB*(i + 1), Bandwidth);
	}



	printf("Benchmarking L2-Cache \n");



	for (int i = 0; i < UsedPointers; i++)
	{
		cudaEventRecord(start);

		BenchMarkCacheKernel << <BlockCount, BlockSize >> >(Pointers[i], 0, Float4Count);

		cudaEventRecord(stop);
		cudaEventSynchronize(stop);

		float milliseconds = 0;
		cudaEventElapsedTime(&milliseconds, start, stop);

		float Bandwidth = (((float)CacheCount* (float)BenchmarkCount * (float)ChunkSize)) / milliseconds / 1000.f / 1000.f;
		printf("L2-Cache-Bandwidth of Chunk no. %i (%i MiByte to %i MiByte):%5.2f GByte/s \n", i, ChunkSizeMB*i, ChunkSizeMB*(i + 1), Bandwidth);
	}


	system("pause");

	cudaDeviceSynchronize();
	cudaDeviceReset();
    return 0;
}


Edit: Ach ja und kennst du eine Möglichkeit die größe des L2$ heraus zu bekommen? Gibt ja keine MSR wie bei einer CPU die das mitteilen
Zum Vergrößern anklicken....
Nur über das Device-Query von CUDA. Ansonsten über Bandbreiten-Latency-Benchmarks.
 
Zuletzt bearbeitet:
Status
Für weitere Antworten geschlossen.
Zurück
Oben

Willkommen auf ComputerBase!

Mit Werbung weiterlesen

Besuche ComputerBase wie gewohnt mit Werbung und Tracking. Die Zustimmung ist jederzeit widerrufbar.

Details im Privacy Center und in der Liste unserer Partner. Ein Widerruf ist möglich in der Datenschutzerklärung.

… oder ComputerBase Pro bestellen

Nutze ComputerBase ohne Werbebanner, Video-Ads und Werbetracking schon für 4 €/Monat oder 36 €/Jahr.

Mehr zu ComputerBase Pro

Bereits Pro-Nutzer? Hier anmelden.

Tracking: Wir und unsere 160 Partner verarbeiten personenbezogene Daten, indem wir mit auf Ihrem Gerät gespeicherten Informationen (z. B. eindeutige Kennungen in Cookies) ein Nutzungsprofil erstellen, um z. B. Anzeigen zu personalisieren. Verarbeitungszwecke: Genaue Standortdaten und Abfrage von Geräteeigenschaften zur Identifikation, Informationen auf einem Gerät speichern und/oder abrufen, Personalisierte Anzeigen und Inhalte, Anzeigen- und Inhaltsmessungen, Erkenntnisse über Zielgruppen und Produktentwicklungen.

Impressum Kontakt Datenschutz