# GPU user experiences

Edit

Thanks to all who have responded! I now have plenty of information on the GeForce 650M GPU as found in most new, mid-range Mac laptops; the Mandelbrot code below runs in comfortably under a tenth of a second. I'm still quite curious how that compares to the GeForce 675MX GPU, which comes in new, high end iMacs. If someone provides that information as an answer, I'll happily accept!

Original

When V8 came out, I was very excited about the introduction of GPU support. So far, however, I've not really been able to take advantage of it the way I would like to, simply because I haven't had a computer with an adequate GPU. Even when CUDA and/or OpenCL are supported, one might not see significant gains when using the GPU due (I think) to limited block size.

Well, now it appears that I'll be getting a new computer in the not so distant future and I'd like to make sure that I'm happy with the GPU. To that end, I wonder if I could receive some feedback on actual user experiences. I am specifically interested in how CUDA and OpenGL work with midline Mac Laptops and Desktops. Laptops seem to run NVIDIA GeForce 640M or 650M GPUs, while desktops are more flexible going up to the 675 or higher in iMacs.

So, does CUDA run on these machines? Specifically, does at least the following return True:

(* Warning - a large package is loaded from Wolfram Research *)
Needs["CUDALink"];
CUDAQ[]

(* Out: False *)


I am aware, of course, of the systems requirements documentation here: http://reference.wolfram.com/mathematica/CUDALink/tutorial/Reference.html#1803279895

That page indicates that CUDA does not run on 600 level GeForce GPUs, but I have a hard time believing that CUDA still is not running on main-line Macs at this point. Furthermore, NVIDIA's information seems to conflict with this here: https://developer.nvidia.com/cuda-gpus

Hence, the question.

Also, I wonder if folks wouldn't mind trying a little test - say, generate a Mandelbrot set, since I'm quite curious about the relative speed of GPUs as accessed through Mathematica. To that end, here's a simple OpenCL program that generates escape times counts to generate a Mandelbrot image. Note that the blockSize parameter can be changed. I'm not an expert on GPU programming, but I believe that higher end GPUs generally allow higher block sizes and that larger block sizes permit more parallelization. My computer allows a blockSize of 16 or lower; lowering the blockSize generally slows down the computation.

Needs["OpenCLLink"];
blockSize = 16;
code = "
__kernel void mandel_kernel(__global Real_t *mSet, int xRes, int yRes,
Real_t xMin, Real_t xMax, Real_t yMin, Real_t yMax) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int i;

Real_t cx = xMin + xIndex*(xMax-xMin)/xRes;
Real_t cy = yMin + yIndex*(yMax-yMin)/yRes;
Real_t x = cx;
Real_t y = cy;
Real_t tmp;

if (xIndex < xRes && yIndex < yRes) {
for (i = 0; i < MAX_ITERATIONS && x*x + y*y <= BOUND_SQUARED; i++) {
tmp = x*x - y*y + cx;
y = 2*x*y + cy;
x = tmp;
}
mSet[xIndex + yIndex*yRes] = i;
}
}
";
If[OpenCLQ[] === True,
mandelCalculate = OpenCLFunctionLoad[code, "mandel_kernel", {{_Real, _, "Output"},
_Integer, _Integer, _Real, _Real, _Real, _Real}, {blockSize, blockSize},
"Defines" -> {"MAX_ITERATIONS" -> 100, "BOUND_SQUARED" -> "4.0"}],
Print["I'm sorry, your computer is even lamer than Mark's!"]
];


Assuming your computer actually passes the test, the following will actually run the program for an xRes by yRes square.

xRes = 1500; yRes = 1500;
(mSet = OpenCLMemoryAllocate[Real, {xRes, yRes}];
mandelCalculate[mSet, xRes, yRes, -2.0, 0.6, -1.3, 1.3];
data = OpenCLMemoryGet[mSet]); // AbsoluteTiming

(* Out: {0.065236, Null} *)


Yeah, that's pretty fast. The computation was performed on the following GPU.

"Renderer" /. ("OnScreen" /. ("OpenGL" /.
SystemInformation["Devices", "GraphicsDevices"]))

(* Out: "ATI Radeon HD 6750M OpenGL Engine" *)


Again, the point is to compare GPUs but, if you want to generate an image, here's one way to do so:

colors = Map[{(100 - #)^2/10000, (100 - #)^3/1000000, (100 - #)/100} &, data, {2}];
Image[colors]

(* Out: Groovy picture *)

-
"Specifically, does at least the following return True" — For a lot of folks, it might not be as simple as that (esp. macs). I know for a fact that on both my MBPs (late 2008 and retina), CUDAQ[] repeatedly gave False even after installing the necessary packages from WRI. In the case of my old laptop, I had to manually install an update from NVIDIA, and in my retina MBP, I had to force the graphics card to always use the NVIDIA card and not the integrated one before I could get them to work (return True). –  The Toad Feb 26 '13 at 22:48
Btw, I get "OpenCLFunction::invblksz: OpenCLLink block size is invalid. " when I run the second code block starting with xRes=... –  The Toad Feb 26 '13 at 22:54
Slightly related : stackoverflow.com/questions/8638905/… –  Artes Feb 26 '13 at 23:07
@rm-rf Per the installation comment - I had a similar experience and needed to install separate software from NVIDIA when trying to get CUDAQ to return True when this stuff first came out. My impression is that this is the reason that Wolfram's data paclets are accessed when first loading CUDALink. –  Mark McClure Feb 26 '13 at 23:08
@MarkMcClure CUDA does run on the 650M, I just ran it this morning. Specifically, you need to download the driver. –  rcollyer Feb 27 '13 at 15:36

I tested your code on my NVIDIA K6000. For some reason, the compiler insisted that I enable double-precision by inserting

# pragma OPENCL EXTENSION cl_khr _fp64: enable


at the beginning of the code (no double quotes used in the code though). For your values xRes = 1500 and yRes = 1500, here are my AbsoluteTiming results:

block_size = 32     {0.031003, Null}

block_size = 16     {0.039004, Null}

block_size =  8     {0.031003, Null}  (a repeat of the block_size = 32 timing value!)

block_size =  4     {0.040004, Null}

block_size =  2     {0.051005, Null}


These timings are not that much faster than yours.

I made the problem one-hundred times bigger, by setting xRes = 15000 and yRes = 15000, and got the following timing results:

block_size = 32     {3.360672, Null}

block_size = 16     {3.364673, Null}

block_size =  8     {3.311662, Null}

block_size =  4     {3.649730, Null}

block_size =  2     {4.830966, Null}


and a second run gave

{4.538908, Null}


To be able to repeat the OpenCL function evaluation for xres = 15000 and yres = 15000 many times, I had to evaluate

OpenCLMemoryUnload[mSet]


after each OpenCL function evaluation in order to recover the large chunk of memory that one evaluation tied up. The time for an evaluation seems to be more or less proportional to the size of the calculation. The largest block_size = 32 seems to give the fastest calculation, but I was surprised by the relative insensitivity of the calculation time to block_size.

The final graphic is indeed groovy, when it is rendered in 2-3 seconds for an xres = 1500 AND yres = 1500 case, but it's not so groovy to wait 2.5 minutes for an xres = 15000 AND yres = 15000 case to render. The big cases do eventually get rendered, but of course they look essentially the same as the smaller cases.

Thanks for posting. I had never run an OpenCL program before. I've only used CUDA.

-