How to max computing throughput?

i think you have confusion between core and thread , you can see this : https://streamcomputing.eu/blog/2017-01-24/many-threads-can-run-gpu/

Thaks for the link but still confused. I see from googling "A9's GPU: Imagination PowerVR GT7600" and choosing Anand link (doing this to avoid URL approval delay) that it should be possible to have thousands of threads scheduled at the same time - but I can only get 512.

# of SIMDs is 6 - this aligns with the "6 GPU cores"

# of MADs (unexplained acronym) 32 - this appears to align with the "threadExecutionWidth" property

But then the total number of MADs is 192, but I can schedule 512 threads total (in 16 thread groups) without seeing any additional delay from just 32 threads. Perhaps my 6s+ has 16 cores?

What I have been hoping is the ability to schedule thousands of threads, not just a few hundred.

I know I'm not thinking about this properly - your article helped - but still don't understand it all.

Thanks so much! And your written English is excellent - 10,000x better than my German :-)

David

Hey David

Glad to hear that. I just noticed that your original question was "how to max computing throughput" and realised that there is another very important thing to know about that. More of a trick exploiting how the things work, and it goes like this: it is often desirable to do more than "one thing" in one thread. For example, if you're doing some kind of image processing, it is only natural to map one image pixel to one thread. And then adjust thread group size/number of thread groups to the image in question. And that will work. But if you do, say, two or four pixels per thread (and change your thread group sizes accordingly), you'll find more often than not that the whole computing will execute faster. This is because there is some constant cost "per thread", part of it induced by how things work, part because how the kernel is written (say, you'll calculate "index" of pixel first, then do actual operations - so every thread has to devote "thread" memory to this "index" variable). If you have same amount of threads, but doing twice the work, costs of processing rise twice (you still have to process same amount of pixels) but these additional costs don't. Depending on scenario, this may give nice speedup. If interested, google "Better performance at lower occupancy" and/or other papers by Vasily Volkov. Unfortunately these are in nVidia CUDA parlance, but general idea applies to many, if not all GPUs. And it certainly worked for me in Metal, too.

Regards

Michal

Again, thanks for the tips and reference! In my case, say calculating the variance of a large array, I can break it up into pieces and do enough work in each thread to essentially have all cores working at once to calculate one value. Suppose I have n number of these - how best to schedule them so when one gets done the other starts up? Should I have multiple CPU threads doing the work, or just loop on these:

// Loop on ...
computeCommandEncoder.endEncoding()
commandBuffer.commit()
commandBuffer.waitUntilCompleted()

This is perfect example how confusing is terminology here. I was refering to Metal "thread", not OS one! So for example something like this:

kernel void variance(device float * array [ [ buffer( 0 ) ] ], constant float mean [ [ buffer( 1 ) ] ],
                     uint2 coords [ [ thread_position_in_grid ] ], uint2 dims [ [ threads_per_grid ] ] )
{
     uint index = coords.y * dims.x + coords.x;
     // do one piece of work per kernel "thread"
     array[ index ] = abs( array[ index ] - mean );     // or whatever you want
}

Is OK, and will work just fine, but you can probably get more performance out of something like this:

kernel void variance2(device float * array [ [ buffer( 0 ) ] ], constant float mean [ [ buffer( 1 ) ] ],
                     uint2 coords [ [ thread_position_in_grid ] ], uint2 dims [ [ threads_per_grid ] ] )
{
     uint index = (coords.y * dims.x + coords.x) * 2;
     // do two pieces of work per kernel "thread"
     array[ index ] = abs( array[ index ] - mean );
     array[ index + 1 ] = abs( array[ index + 1 ] - mean );
}

Now you run the second version with X resolution halved, so you have half the execution threads (again, Metal "threads", not CPU ones!). It should

give you same results (assuming that you can halve X resolution and all). But it would probably be faster. Hope it is clear now.

Michal

PS. Please understand that these are just sketches, not actual working code. I just wrote them there to illustrate general idea.

PPS. As for CPU-side Metal threading, I know that it is possible (there are threaded command encoders) but I never actually used them, and from what I understand these are for special cases, like rendering a lot of objects with different properties. Something that will happen more often in computer game area, I guess, less in actual numerical computations.