У меня есть один трёхмерный блок (или несколько, но это не важно, так как они независимы). Мне надо синхронизировать потоки в этом блоке следующим образом : сначала должны исполнится все потоки с координатой z=0, затем z=1 и так далее. То есть надо послойно синхронизировать потоки по оси z. Мне уже несколько человек приводят в пример следующий код:
__global__ void Kernel(){
for(int z = 0; z < zmax; z++){
if(threadIdx.z == z){
//code
__syncthreads();
}
}
}
Если для наглядности сократить это всё, то вместо цикла будет что-то вроде этого
if(threadIdx.z == 0){
__syncthreads();
}
if(threadIdx.z == 1){
__syncthreads();
}
и т. д.
Такой код по идее гарантирует синхронное выполнение потоков с одинаковым z, но не гарантирует порядок z. Ведь могут первее к точке синхронизации придти все потоки с z=1, и тогда слой, где z=1 выполнится раньше того, где z=0.
Собственно верны ли мои рассуждения? Если да, то как сделать так как я хочу. Или это невозможно?