StackOverflow.com, under the "cuda" tag, would be a good place to ask questions like this. Or devtalk.nvidia.com (developer forums).
It is not defined as such by nvcc or the cuda headers, but for all intents and purposes it can be made one, right (while giving up some level of portability)?
If it is a concern that code compiled now may run e.g. in a decade from now on a 64-wide warp arch, one can assert(WARP_SIZE == warpSize) for safety.
In the sample you have
deviceReduceKernel<<<blocks, threads="">>>(in, out, N);
what does the ' threads="" ' mean in the code above?
It was a parsing error. I've fixed it. Thanks!
You said:
"If the source lane ID is out of range or the source thread has exited, the calling threadâs own var is returned."
I encountered a case where some arbitrary value instead of the thread's own one was returned by __shfl_down accessing a variable of a thread which wasn't started by my kernel.
I adapted your code for a parallel 3D AABB computation.
It didn't work when the block dimension was 95 and the warp size 32.
The problem is: blockDim % warp size != 0.
I never started the 96th thread which resulted in a wrong reduction result in the function corresponding to warpAllReduceSum for the last warp.
In my case 0 was returned by __shfl_down for the x-axis although all x-coordinates were negative.
This leads to the following question:
When do I need to take care of the shuffle operations accessing invalid threads?
Is it valid if a lane ID outside the warp is computed?
Is it invalid if a lane ID inside the current warp is computed for a thread that wasn't actually started?
Are there other cases to be considered?
According to the CUDA C Programming Guide:
Threads may only read data from another thread which is actively participating in the __shfl() command. If the target thread is inactive, the retrieved value is undefined.
http://docs.nvidia.com/cuda...
I think this answers both of your cases, as they are two instances of the same thing: if the lane ID is outside the warp, retrieved value is undefined. If the lane ID refers to an inactive thread, retrieved value is undefined.
__shfl() is misspelled as __shlf()
Justin, thanks, it is very helpful info!
Please note following line is wrong:
val = (threadIdx.x < blockDim.x / warpSize) ? shared[lane] : 0;
It should be:
val = (threadIdx.x <= (blockDim.x-1) / warpSize) ? shared[lane] : 0;
Hi Vadim. Actually, this line of code is correct as written. blockDim.x / warpSize gives the number of warps, and since threadIdx.x is zero-based, using <= would cause an out of bounds error.
Mark,
thx for quick reply. But I disagree, just an example: what if blockDim.x = warpSize+1 ? Say for warpSize=8:
>>> [int(x<9/8) for x in range(8)]
[1, 0, 0, 0, 0, 0, 0, 0]
So, second warp's sum is nullified.
But correct code would be:
>>> [int(x<=(9-1)/8) for x in range(8)]
[1, 1, 0, 0, 0, 0, 0, 0]
Yes, the code assumes the block size is a multiple of the warp size. While the code is not bulletproof because of this, it's a common assumption.
Ok, but this simple change allows your block to be not a multiple of the warp. Does it make sense?
Thanks Vadim, let me clarify a bit. The post assumed the block size is divisible by the warpSize. This has to be true in order to use __shfl safely. __shfl on a partial warp is considered undefined. You may get correct results but you are not guaranteed to. We will update the post to clarify this.
I am getting always correct results for partially filled warps (with inactive lanes), and checked all different combinations on 3.0 and 3.5 capable devices, also shared by many processes. So I wonder what condition triggers registers to be non zero initialized. Is it something tunable in the driver maybe?
Nope. It is undefined. Undefined means that we reserve the right to change the behavior at any time (for example, new driver, new hardware, new toolkit, etc).
It would be amazing if the whole code was global, I mean that can work in Fermi and Kepler architecture.
Hi, this is a great post! Thanks for sharing it.
Is it possible to have a kernel, where I fill the vector and after that I can call
your reduce function, instead of copying the vector to Host and Copy it back to the reduce function? It would save time by doing that.
Thanks.
sorry but i can't understand this: in BLOCK REDUCE
i use a block with 1024 threads; each warp calculates its own reduction within 32 threads; 1°question: what value has val? 2°question: the meaning of
val = (threadIdx.x < blockDim.x / warpSize) ? shared[lane] : 0;
Justin, Very interesting. In the case using atomics you have to reset "out" on before every reduction; what is the recommended way to do this since I understand that out must live in device memory?
excellent article! took me a little to digest it but it is super smart. thank you for taking time to write this, the initial loop grid-wide is especially smart! also because it naturally handles boundary conditions.