Yeah seems like maybe defect in cache RAM… and then that one gets disabled… rest re-routed or placed differently on chip. Or maybe it was disabled to bring down heat ? Wwoow… Or maybe just a human mistake… an overlooked bottleneck problem.
Hello,
I stumbled upon this obscure (for a non-developer…) part of the site, and I am willing to help. I do not know, however, what parameters I should use when I am running your app. So I innovated a bit for the first run :)
Memory Block Size: 134217728 (default)
Rounds: 2048
5 iterations at default settings (3 rounds)
About heat: no, the card performs exemplary, at least with the Gainward custom cooling. It never goes past 81 Celsius (target is 79 as default). It has lots of headroom and can easily be overclocked. Its boost goes to 1392 MHz in my case most of the time even the specs say 1304 MHz.
There is a HUGE drama-scandal going on at the Geforce forums and even I have decided to RMA, but I suppose this is another issue.
All in all, I might not be a CUDA dev, I am just a gamer/graphic designer, but if you can discover any other details about the VRAM capabilities, I will help, as I see nobody offered.
I will not make any details about your app public at the Geforce forums, if you don’t want to, but you would get tens and tens of voluntaries if you decide to share it, as this place is definitely off the map and I am wondering myself how did I found it on Google.
Let me know if I can help further, run at different parameters etc.
Thanks for taking the time to write this app.
I tested my 970. The result seem to differ from undone’s. My results match more the original Nai’s benchmark results.
[url]http://imgur.com/bKV5mmp[/url]
I managed to reproduce your results by putting the card on Prefer Maximum Performance and using the High Performance power mode. Are these your settings?
Here is how it looks.
3x1024 runs:
10x3 runs:
Interestingly enough, my videocard somehow interfering with my old audio card/speakers, and M-Audio Audiophile USB/Yamaha HS80Ms. If I run the larger 1024 thing, you can hear the GPU ( I think…) buzzing in the speakers… but when it reaches that huge drop, all noises stop, as if nothing even happens. So strange.
There are already people claiming that those 512 MB slow VRAM are not even used, the card uses normal RAM/pagefile instead.
Oh well.
Extremely odd. After capturing these 2 images, I was completely unable to reproduce the “nai” behavior again, and started getting 110-150 GB erratic bandwidth again. As of now, I just can’t reproduce the 20 GB/s drops at the last memory blocks. procexp never shows the videocard using more than 3 GB VRAM too…
Found something worrying, if something can worry me about this card anymore…
Did’t observe up to now.
Both mine and briose’s screenshots show 22 (23?) blocks tested. Well that might be incorrect.
The benchmark itself reports:
Memory Total: 4294967295
Memory Free: 3794739200
Memory Blocks:28 (29 reported if I empty the most VRAM I can, with basic scheme etc.)
Looking at Nai’s benchmark, you can see it allocates 29-30 “chunks”. People recommended that you run the benchmark after booting from your iGPU, but I don’t have any. Lowest VRAM utilization I managed to get within Windows 8.1 is like 100 MB.
I assume that the “missing” memory blocks are filled and untested. It’s kinda interested that it looks like Windows 8.1 and the driver chose to put the OS files in the slow 0.5 GB from what I can make out of the tests.
I think it is very important that regardless of having 28 or 29 blocks available, the card will ALWAYS test just 22 and always fill just 3 GB VRAM. Looks quite obvious to me the driver/BIOS is doing its best to never allow anything probe the slow memory zone.
Wish Skybuck was here to help us interpret the results, and hopefully agree to sharing this app, so we get more tests coming in. The guys at the Geforce forums would be more than happy to test it. I didn’t say a word about it yet.
After starting PC, I tested and graph did not show the dip at the end. It was fluctuating between 130-140 all the time. Then I started Dying Light for a minute and ran the test again. Now there is the dip again at the end.
Also the logs show that I have 27 memory blocks. Might be due to running 2 monitors currently.
My power management mode is set to adaptive. And the windows power plan is set to balanced.
Very strange. It looks like the app cannot test some areas, also that the card tries to avoid access to the low speed zone. It’s all that I can make out of it, would be great if the developer would return to take a look at the results.
Pretty much a 3.5 GB VRAM card for all intents and purposes.
Definitely, at least with my tests in Shadow of Mordor and Far Cry 4, there are lots of random stutters when GPU mem usage reaches 3600MB. Especially Shadow of Mordor is butter smooth 100FPS when mem usage is around 3200MB. I really should test a bit more though.
And I bought two 970s. I figured for once I would get the killer setup for several years… “#%&”#¤!!!
Hi, thank you for trying out my app. The app allocates blocks according to free memory available. I currently have no explanation why the app only tests 23 blocks. Perhaps if you send the log to mail e-mail address at skybuck2000@hotmail.com I could take a further look at it.
There is also something else I’d be interested in seeing. For example try this setting:
MemoryBlockSize: 1048576
(This will test the graphics card in chunks of 1 MB instead of 128 MB.) If you do decide to perform that test please post screenshots of the graphs here. That would be most interesting to look at.
What’s also somewhat remarkable is the steep drop at the last couple of blocks, though it seems to go high again at the very last one. This could be because the RAM might be virtualized, so it might be any block that could drop in performance, I am not sure about that.
Also you can see that you can always test about 3 GB. procexp shows VRAM being filled with 3GB+what was previously in the VRAM.
Logs give CUDA error when trying to access those last memory blocks. Only noticed after you asked for them and I scrolled down to see if there’s more text :)
Memory Block Number: 24
Error: Memory Block Allocation Failed: TCudaDriverMemory.SetSize: AllocateMemory( 134217728 ) failed.
Memory Block Number: 25
Error: Memory Block Allocation Failed: TCudaDriverMemory.SetSize: AllocateMemory( 134217728 ) failed.
Memory Block Number: 26
Error: Memory Block Allocation Failed: TCudaDriverMemory.SetSize: AllocateMemory( 134217728 ) failed.
Memory Block Number: 27
Error: Memory Block Allocation Failed: TCudaDriverMemory.SetSize: AllocateMemory( 134217728 ) failed.
Memory Block Number: 28
Error: Memory Block Allocation Failed: TCudaDriverMemory.SetSize: AllocateMemory( 134217728 ) failed.
Some notes and things I may try in the future:
-
First of all the app doesn’t show maximum bandwidth accurately yet. The test completes too fast on fast cards, or too fast with small memory block settings. This is because the accuracy of the cuda events is limited to milliseconds, the test will complete in 1 millisecond or less, making it impossible to measure the true time it takes. One solution could be to adjust the kernel so it loops multiple times and incorporate this in the bandwidth calculation. Another thing I am interested in is estimating the bandwidth a device can have based on attributes/device query/information, this will be more difficult to do because of architectural differences between gpu cards, these architectural differences are not mentioned in this kind of query return information, for example number of store/load units could be different but is not mentioned. However there is some basic information like memory clock frequency, bus width, as far as I know, double/tripple memory rate also not specificed so it will get a bit tricky/inaccurate prediction, but perhaps it will be good enough to give the kernel some kind of indication how much bandwidth it can except and adjust it’s loops accordingly also based on memory block size, alternatively I could also let the user specify expected bandwidth rate and go from there, there might be more easy for a first version.
-
It’s also possible to optimize the kernel/ptx further by reducing unnecessary instructions in case the test/kernel/device is limited by processing speed:
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sat Mar 15 02:52:22 2014 (1394848342)
// Cuda compilation tools, release 6.0, V6.0.1
//
.version 4.0
.target sm_20
.address_size 32
.visible .entry BenchMarkDRAMKernel(
.param .u32 BenchMarkDRAMKernel_param_0
)
{
.reg .pred %p<2>;
.reg .s32 %r<9>;
.reg .f32 %f<18>;
ld.param.u32 %r2, [BenchMarkDRAMKernel_param_0];
cvta.to.global.u32 %r1, %r2;
mov.u32 %r3, %ctaid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %tid.x;
mad.lo.s32 %r6, %r3, %r4, %r5;
shl.b32 %r7, %r6, 4;
add.s32 %r8, %r1, %r7;
ld.global.v4.f32 {%f5, %f6, %f7, %f8}, [%r8];
BB0_2:
ret;
}
(Stuff that in the *.ptx file).
-
I may also try a multi-threaded version… just to make gui more responsive, maybe multi threaded version will also give more consistent performance, or perhaps worse, time will tell ;).
-
I already tried an overhead graph to see if that could explain why some mentioned some blocks not allocated, my hypotheses was perhaps there is huge overhead associated with allocating multiple blocks, this does not seem to be the case on my fermi/gt 520 card, however there was some weird overhead at the last view blocks, as if the memory manager got confused and allocated the same block twice or something or re-used a previously allocated block, a bit weird. So far overhead seems to be 0 most of the time.
-
One theory/hypotheses for later memory blocks not to be allocated is “memory fragmentation” perhaps some other process is still holding on to some variables inside the graphics memory. Thus seems a bit weird ?! Should those page blocks not be swapped to some place else ? Hmm kinda strange. (Assuming a virtual memory design, that should be possible)
-
I also read the device memory/driver manages it via a linked list, but why this would reduce bandwidth for a 4 MB memory block test (recommended for all to try) I don’t understand.
On my GT 520, the performance goes then down somewhat.
-
So I may also try and allocate as much memory as possible and test those regions inside that big memory chunk just to see if this matters or not this will be a secondary test/allocation mode.
-
Perhaps more information in log and a button to save it directly to file or/and perhaps a checkbox to save it directly to file as the test progresses in case the app crashes or anything like that so at least some parts could be captured to a file.
-
Correction of chart/memory ammount displayed, it should be expanded to X.X GB to show fragment of GB.
-
Re-compile of kernel with cuda 6.5 or cuda 7 or something like that to see if it matters or not or to see what happens.
-
Perhaps examine the run in debug driver with nvidia profile or nvidia charts/insight to see what happens further ;) to see what bottlenecks could be or what insight recommends.
-
Possibly re-write/write/change to new self-developed kernel instead of Nai’s kernel so that kernel is a bit more parallel, perhaps temp variable a bottleneck, also perhaps fill memory with random values, use C’s “overflow allowed” feature to checksum all values and check the final value (though this could again introduce a bottleneck if all values to be summed together) so perhaps ignore checksum idea, and stay with manually modifieing kernel idea to reduce unnecessary instructions, would be nice if somehow it could be verified that the kernal is indeed reading memory like it should without actually having to process the memory. To prevent processing bottleneck. However testing a minimum ammount of processing could also be interesting to test that. How to do parallel processing and parallel verification of processing still needs to be invented ;) perhaps parallel checksums or something. Or perhaps a checksum per thread, that could be a nice idea.
These are my ideas for further work on the app and further investigation. So stay tuned ! ;)
Skybuck’s CUDA Memory Bandwidth Performance Test version 0.05 is now available:
http://www.skybuck.org/CUDA/BandwidthTest/version%200.05/
Most important changes:
-
Estimated bandwidth based on memory clock frequency and bus bit width and double data rate assumption. (User can override this setting, for example higher to increase load, or lower for faster testing, currently my estimation is that current estimation will put 99% load on gpu !).
-
Ammount of threads “spawned” is now based on estimated bandwidth, thread index is wrapped around memory block size (items that will fit inside of it).
No further features implemented yet from list.
Additional feature for future:
12. update compute capabilities/specifications for more accurate optimal dimension calculations for newer graphics cards. (Calculations already updated to 64 bit to be able to do this program).
In short: This app should now display bandwidth much more accurately. Note: Lower memory block settings may give slightly higher bandwidth numbers because of cache effects.
Alright, what should we fill in this time for Memory Block Size (134217728 default) and Estimated Bandwidth (-50557906944 default, yes negative…).
nVIDIA still advertises Memory Bandwidth (GB/sec): 224, but we now know this ain’t true, with the 3.5GB being much faster etc.
This is the Geforce forums topic:
https://forums.geforce.com/default/topic/809501/geforce-900-series/skybucks-vram-cuda-bandwidth-performance-test
Fell free to drop by and maybe offer some support if people are lost. Keep in mind most users will be gamers not coding experts.
I tried to run at 224000000000 bytes btw and I didn’t get any meaningful graph.
Log:
i got it to test upto block 11888 but i clicked on the settings tab to note the settings i used but the program stopped responding for a while (middle pic) then stopped (bottom pic) took these screen shots if they help
[/URL][/img]
i have to turn pc off and reboot pc to get est to runn again you also must leave your PC alone or it stops the test,thats what i have found
i will re-run it and see what happens
these are my settings
mem block size = 1048576
rounds = 52
estimated bandwidth = 2240
SLI disabled
[/URL][/img]
another run with settings
[/URL][/img]
maybe this image will help tell you more test run in SLI GTX970’s Game TitanFall max settings, monitor by Nvidia Inspector 1.9.7.3, Game Rez run in DSR 3840x2160
[/URL][/img]
FarCry 3, Max Game settings, DSR 3840x2160 Game Rezz, Monitor GPU-Z 0.8.1
[/URL][/img]
Note the GPU LOAD, Memory Controller And Bus Interface Load all go crazy once ram goes above 3.6G vram Usage.
OK this is my last test for you maybe now you have at least a bit more info to look at
better image here
[/URL][/img]
Hi,
Thank you for testing version 0.05. The estimated bandwidth in this version would indeed go negative, however if the user manually corrected it then it might still work funny enough.
However:
Skybuck’s CUDA Memory Bandwidth Performance Test version 0.06 is now available !
Big improvements in this new version:
- Multi-threaded support (it’s now dual threaded, so GUI remains responsive).
- ComboBox to select memory block size more easily.
- Negative estimated bandwidth fixed.
- Much more info in Log.
- Save log to file button.
- Start/Stop functionality to stop test prematurely.
- Compute capability 5.x supported.
I would recommend to at least run with a memory block size of 1 MB. I would also recommend to first perform a small test with 128 MB block, and then measure actual bandwidth. Then use the measured number and put it into estimated bandwidth to overwrite it. This should reduce test time closer to 1 second per memory block. My GT 520 is supposed to be able to achieve 9 GB/sec or so but it only achieve somewhere between 2.7 GB/sec and 3.5 GB/sec. Not sure yet why it cannot achieve higher, maybe compute-bound.
Anyway here is the new version:
http://www.skybuck.org/CUDA/BandwidthTest/version%200.06/
I have also uploaded some new charts of my GT 520 (and logs):
http://www.skybuck.org/CUDA/BandwidthTest/Charts/
(Perhaps later I will upload a 1 MB memory block chart)
I think the driver might have some bandwidth/compute overhead as soon as the number of blocks increases, the performance would go down. This is just a hypothesis which I like to explore in a future version which will support allocating just 1 big block and then processing that big block in chunks, just like this version does except it won’t allocate many chunks, just 1 big one.