NVIDIA GF100 & Friends speculation
- Thread starter Arty
- Start date
Why would you need that, isn't a pointer just an index into global memory?(or some UAV in this case). And with the interlocked counters you'll have no problem "allocating" new blocks either. I see no problem in either building or traversing such structures.
Meh, yeah I guess for something like an octree you could use an append buffer to add child nodes and just update the parent with the index for each child. Same effect as a pointer but ugly
Btw, do GPU's currently support mutexes? That would be a requirement for rebuilding the acceleration structure on the GPU no? (say two threads want to create the same child node in a kd-tree to insert two different objects that fall into that cell)
aaronspink
Veteran
Meh, yeah I guess for something like an octree you could use an append buffer to add child nodes and just update the parent with the index for each child. Same effect as a pointer but ugly
Btw, do GPU's currently support mutexes? That would be a requirement for rebuilding the acceleration structure on the GPU no? (say two threads want to create the same child node in a kd-tree to insert two different objects that fall into that cell)
I don't believe they currently support the cpu style light weight mutexes such as LL/SC or FetchAndAdd/CmpExchange but they do support sync primitives that could be used as a heavy weight alternative. But in general I think you are currently better off not playing around with any link list style structures (LL, n-ary trees, latices, etc).
CUDA supports atomicCAS(), but it is relatively slow at the moment. Building mutexes out of it is not a particularly good idea on the current generation.
Actually, I don't think you'd want to use a syncthreads_count() to do a prefix sum on the mask. syncthreads_count() only returns the reduction of all the predicates in the thread block, which is not the same thing at all.
Ballot is actually very useful for prefix sums, though: you can use it in conjunction with a boolean mask and a popcount to do a warp scan efficiently, for example:
uint warpId = threadIdx.x & 0x1f;
uint warpMask = (1 << warpId) - 1;
...
int predicate = foo();
uint ballot = __ballot();
uint parents = ballot & warpMask;
int warpScan = __popc(parents);
Interesting, nice workaround
Cool, I wasn't aware that the __popc function existed. So yeah, in conjunction with __ballot you get a warp level sum. I still don't understand why you think syncthreads_count isn't useful though. There will be times when you want to evaluate the mask over the entire block. Essentially it's doing everything above as well as doing a final sum of all the individual warpScan values for thread(warpsize) across all warps.
If you mean that syncthreads_count isn't useful as a scan on a single warp level mask produced by __ballot then yeah, absolutely. I should've worded my original comments better to make it clear that I wasn't suggesting otherwise.
Jawed
Legend
http://forum.beyond3d.com/showthread.php?p=1327227#post1327227
For the correctly linked patents being referred to.
Yes, the description of __ballot() doesn't match with the PSCAN instruction. I'm confused why PSCAN in that code snippet is not in the the CUDA 3.0 guide
Is it missing?
Jawed
RecessionCone's macro seems to be the closest thing to it. There's no single instruction that accomplishes what PSCAN does. Also, I was wrong about syncthreads_count as it returns a single value to all threads - the count of true predicate evaluations. At first I thought it was similiar to rank() from the patent but meh.
This would be a real fun because i read somewhere that GT200b is EOL.
So because of their Quadro and Tesla line there is now a surplus of GTX200 series gpus floating about on newegg? So close to their next gen release? Bah, I don't get it myself...how can there be no stock for months then all of sudden, close to their release of next gen there is stock? I get the impression this will be their mid-range gpus...
I do think syncthreads_count is useful, but I'm trying to point out that it does something very different: it's a reduction, rather than a prefix sum. Just to be explicit:
x = [1 2 3 4 5]
reduce(x) = 15 # You can get this with syncthreads_count
prefix_sum(x) = [1 3 6 10 15] #You can get this with ballot, etc.
You can derive all the elements in a prefix sum across a warp using ballot, as I pointed out. You can't derive all the elements in a prefix sum either across a warp or across a thread block with syncthreads_count. It produces a single scalar, the reduction of all the predicates from all threads in the block, after the block synchronized. This is much less information than you can get with a proper prefix sum.
For example, you can use a predicate prefix sum to compact results, if every thread may be producing 0 or 1 elements, you can use the prefix sum to calculate the address to which each particular thread should store their result. The result of syncthreads_count, while still useful, only gives you the total amount of results the block will be producing. They're fundamentally different things.
In any case, I'm happy about both of them. =) You might as well do some computation as you hit the barrier, so I'm happy they're providing any and all and count barriers. But that's very different from a prefix sum. :smile:
Yep, thanks. I realized my mistake a little while ago
. syncthreads_count was appropriately named aaronspink
Veteran
It's prudent for any company to care what their competition will do.
What would that prove other than showing it can't retire kernels out of order? There can still be 16 running in parallel even if the chip hangs.
Similar threads
