order of execution in a divergent warp

Accelerated Computing CUDA CUDA Programming and Performance
#21 
for(int i=threadIdx.x; i<SHARED_MEM_CAPACITY; i += blockDim.x)

{

  sharedArray[i] = INIT_VALUE;

}

__syncthreads();

Wont this initialize a shared array with less than 32 threads a block?

#22

I was replying to the initial post… but not to your post on re-convergence.

#23

I was replying to the initial post… but not to your post on re-convergence.

#24

Yes, but that still requires reconvergence before the __syncthreads() if SHARED_MEM_CAPACITY is not multiple of warp size.

#25

Yes, but that still requires reconvergence before the __syncthreads() if SHARED_MEM_CAPACITY is not multiple of warp size.

#26

I would be surprised if the code above does not work anymore. May be, their compiler has become too smart.

#27

I would be surprised if the code above does not work anymore. May be, their compiler has become too smart.

#28

Well we were suprised too :D
But the code you posted is probably safer as it is harder for the compiler to optimize + there is __syncthreads.