Basically I have two large data arrays (226799 uint4) and the algorithm can be split into two sub algorithms, so each half-iteration processes one sub algorithm, let’s call them A and B, and data arrays 0 and 1.
A reads from 0 performs computing on the values read and writes to 1.
B reads from 1 performs computing on the values read and writes to 0.
Each thread reads aligned data (data is accessed contiguously) but writes in a “sparse” way to distinct memory locations.
As the amount of data is too big, no grid that can be active at the same time in the GPU is possible, so no inter-block sync there.
So for a given number of iterations, the kernel must be called from host that same amount of times. If the loops is included on the kernel side, write-after-read hazards occur, only for data accessed by the blocks that haven’t finished their work on A and B is already reading from them. The contiguous reading and sparse writing could be turned into a sparse reading and contiguous writing but the problems would be the same. That said, I can’t allocate data in shared and compute the iterations I want and then merge it on global memory.
Basically I have two large data arrays (226799 uint4) and the algorithm can be split into two sub algorithms, so each half-iteration processes one sub algorithm, let’s call them A and B, and data arrays 0 and 1.
A reads from 0 performs computing on the values read and writes to 1.
B reads from 1 performs computing on the values read and writes to 0.
Each thread reads aligned data (data is accessed contiguously) but writes in a “sparse” way to distinct memory locations.
As the amount of data is too big, no grid that can be active at the same time in the GPU is possible, so no inter-block sync there.
So for a given number of iterations, the kernel must be called from host that same amount of times. If the loops is included on the kernel side, write-after-read hazards occur, only for data accessed by the blocks that haven’t finished their work on A and B is already reading from them. The contiguous reading and sparse writing could be turned into a sparse reading and contiguous writing but the problems would be the same. That said, I can’t allocate data in shared and compute the iterations I want and then merge it on global memory.
I think my solution will work. It should work for two steps too. How many iterations have you? “Data size” in my code means not pure data size in bytes.
“As the amount of data is too big, no grid that can be active at the same time in the GPU is possible, so no inter-block sync there.” It is a bit wrong.
I suppose you need to convert many frames. “So for a given number of iterations, the kernel must be called from host that same amount of times.” It is not true too.
Maybe you have some misunderstanding, when you set 1000 blocks, they are not executed on GPU at once simultaneously.
I think my solution will work. It should work for two steps too. How many iterations have you? “Data size” in my code means not pure data size in bytes.
“As the amount of data is too big, no grid that can be active at the same time in the GPU is possible, so no inter-block sync there.” It is a bit wrong.
I suppose you need to convert many frames. “So for a given number of iterations, the kernel must be called from host that same amount of times.” It is not true too.
Maybe you have some misunderstanding, when you set 1000 blocks, they are not executed on GPU at once simultaneously.
__device__ void kernelA(/*ptrs*)
__device__ void kernelB(/*ptrs*)
__global__ void iterative_kernel(/*ptrs*)
{
for(/*given number of iterations*/)
{
kernelA
kernelB
}
}
As threads in GPU are created and then occupy the SM until their death, having a grid that can’t be fully loaded on GPU implies that no dependency from block X that isn’t loaded can be possible, as all other blocks occupy the resources the block X needs to be executed so deadlock is ensured.
If __interblocksync() exists:
__device__ void kernelA(/*ptrs*)
__device__ void kernelB(/*ptrs*)
__global__ void iterative_kernel(/*ptrs*)
{
for(/*given number of iterations*/)
{
kernelA
__interblocksync();
kernelB
}
}
This would imply grid fully loadable on device to avoid deadlock. But that would give a for loop inside the kernelA and kernelB to make each of their threads to access more data than they do. Besides, compiler can’t cope with the too many lines of code a call to a loop with many iterations as functions are inline. With two iterations I get that warning of too many lines in code, compiler not optimizing et cetera…
Your solutions works for 500 blocks whilst only 300 can be running?
CUDA Programming Guide example on data producer data consumer works, but at the end only one block is active…
__device__ void kernelA(/*ptrs*)
__device__ void kernelB(/*ptrs*)
__global__ void iterative_kernel(/*ptrs*)
{
for(/*given number of iterations*/)
{
kernelA
kernelB
}
}
As threads in GPU are created and then occupy the SM until their death, having a grid that can’t be fully loaded on GPU implies that no dependency from block X that isn’t loaded can be possible, as all other blocks occupy the resources the block X needs to be executed so deadlock is ensured.
If __interblocksync() exists:
__device__ void kernelA(/*ptrs*)
__device__ void kernelB(/*ptrs*)
__global__ void iterative_kernel(/*ptrs*)
{
for(/*given number of iterations*/)
{
kernelA
__interblocksync();
kernelB
}
}
This would imply grid fully loadable on device to avoid deadlock. But that would give a for loop inside the kernelA and kernelB to make each of their threads to access more data than they do. Besides, compiler can’t cope with the too many lines of code a call to a loop with many iterations as functions are inline. With two iterations I get that warning of too many lines in code, compiler not optimizing et cetera…
Your solutions works for 500 blocks whilst only 300 can be running?
CUDA Programming Guide example on data producer data consumer works, but at the end only one block is active…
GPU could perform about 30*8=240 blocks max at one time. So new blocks replace completed blocks on gpu. Btw, it is not important here. Block just need to know its start number. If it is less than for example 100, it performs first stage, if it is more than 100, it performs second stage. And after start it needs to wait all blocks from first iterations to be completed. It is simple, block increments global block counter at block’s end.
GPU could perform about 30*8=240 blocks max at one time. So new blocks replace completed blocks on gpu. Btw, it is not important here. Block just need to know its start number. If it is less than for example 100, it performs first stage, if it is more than 100, it performs second stage. And after start it needs to wait all blocks from first iterations to be completed. It is simple, block increments global block counter at block’s end.
Suppose 308+1 blocks need to run, and 308 are running, those blocks will eventually end up in a deadlock as block 30*8+1 th can’t run and the others blocks will eventually wait for that block. It’s not like a host thread that can have at a given time 0% CPU time controlled by programmer. Threads are spawned and run until they die, if they have to wait for something, they’ll wait, whether if that in the whole grid leads to deadlock.
Suppose 308+1 blocks need to run, and 308 are running, those blocks will eventually end up in a deadlock as block 30*8+1 th can’t run and the others blocks will eventually wait for that block. It’s not like a host thread that can have at a given time 0% CPU time controlled by programmer. Threads are spawned and run until they die, if they have to wait for something, they’ll wait, whether if that in the whole grid leads to deadlock.
Why those blocks will wait for block 30*8+1? My code will not deadlock for sure. Only problem is to place __threadfence in right place. But I know now where it should be placed. If it works like it was described. Only problem is performance, but it should be faster than kernel launch.
I only missed __threadfence() at the end of the blocks, I confused it with __threadfence_system().
Why those blocks will wait for block 30*8+1? My code will not deadlock for sure. Only problem is to place __threadfence in right place. But I know now where it should be placed. If it works like it was described. Only problem is performance, but it should be faster than kernel launch.
I only missed __threadfence() at the end of the blocks, I confused it with __threadfence_system().
I was looking for a global barrier on a GPU environment, and found this topic. I’m a newbie on GPU programming, so sorry if my question doesn’t make sense, but: What is the purpose of an barrier that doesn’t guarantee memory consistency? What kind of application is it useful?
I was confused about the various fence() functions for a while, but I’ve come to think of them as flush commands for the memory hierarchy. Fermi devices do not have a coherent cache (unlike most multi-core CPUs), so it is possible that the most up-to-date value for a global memory location could be found at different places:
a thread register (if a global write is deferred to later by the compiler)
L1 cache
L2 cache
device memory or host memory (the latter is possible if you have mapped host memory into the GPU memory space)
At the end of a kernel, all writes are guaranteed to be flushed all the way up the hierarchy to device memory (or host memory). During kernel execution, nothing is guaranteed beyond internal thread consistency unless you explicitly ask for it using a fence function:
__threadfence_block() flushes writes up to the L1 level, so other threads in the block can see the write.
__threadfence() flushes writes up to the L2 level, so any thread on the device can see the write
__threadfence_system() flushes all the way up to host memory, so the host CPU can see the write
(As a bonus, all the fence functions also flush writes to shared memory, which is also one of the two jobs of syncthreads().)
You might want to use a memory fence function to implement non-blocking communication between threads. The CUDA Programming Guide gives the example of using atomic operations (which go directly to L2) to set a counter, indicating that another block can read some output. In this case, __threadfence() is called first to ensure the data has been flushed up to the L2 cache before incrementing the atomic counter. Without __threadfence(), it is possible for the consumer of the data to get stale values from the cache.
Hopefully that helps clarify things a bit. I would consider memory fencing an advanced GPU programming technique that is important when you need it, but most of the time you don’t need it. It is much easier to restrict yourself to shared memory and __syncthreads() for intra-block communication and barriers, atomics for limited inter-block communication, and queuing multiple kernels to create global thread barriers between kernels.
Thanks, it is more clear now. I’ve never used fences, but I’ve developed many parallel programs, so I have an idea about what it means. Based on previous replies, fences are slow, so “don’t use it, unless you really need”. It seems that for majority of the situations, multiple calls for the kernel is easier and it will have almost the same performance.
Another point is (correct me if I’m wrong), but it is possible to create N blocks, but there is no guarantee that all of them will be active at the same time. Some of them will be queued waiting the for the active blocks to finish their execution. In this scenario, a global barrier doesn’t seem to make sense.
Well, possibly you can find a way around. First, you can try to pick such number of blocks that will happen to be active at the same time. Second, you can try scheming a voting algorithm to make the currently active blocks find each other and synchronize only them. Etc.
I’d like to point out though that these are rather advanced programming techniques. While I use them in my own code (and they work well), I’d rather not want to guide someone to use them through the forum who is just picking up CUDA.
I am looking into developing a global barrier for a real-time algorithm, with concurrent kernels, very few blocks each. The motivation is to batch together multiple kernel invocations into concurrent kernels that run longer and synchronize on the device, without getting back to the host, ultimately reducing the overhead of global synchronization through the host. This thread seems related and so I thought I’d ask a few questions.
Any pointers regarding concurrent kernel execution and global barrier implementations?
Do we have any guarantees w.r.t. the concurrent kernel block launches. Suppose one kernel has only 1 block, and another has 16. Would they ever run concurrently? E.g. the first kernel takes 1 SM, and let’s say 14 more SMs are available. Would the device run-time schedule the second kernel to run concurrently? I assume that’s the case.
How about the other way around, would the run-time launch 14 blocks from the second kernel and 1 block from the first kernel when both kernels are available “at the same time”?
Suppose I launch two kernels with just one block each, over two independent streams. These kernels may want to run concurrently for a “long” time, and synchronize through a global barrier, working in a pipeline fashion, passing memory buffers to each other. Is there any reason the device run-time would want to run one kernel to completion before even starting the other? Again, each kernel has just one block, so each would take just one SM.
There are no guarantees at all. So you should definitely test whether the different kernel’s block “see” each other. And fail the whole invocation (and fall back to a safe implementation) when they don’t.