The concept of a warp is relevant for this second kernel, as sequential memory accesses by threads that are part of the same warp can be grouped and executed as one.
is it because they are executed in a lockstep ?
Also where is the data stored after coalescing
TODO: GFLOPS and its calculation
__global__ void staticReverse(int *d, int n)
{
__shared__ int s[64];
int t = threadIdx.x;
int tr = n-t-1;
s[t] = d[t];
__syncthreads();
d[t] = s[tr];
}If the kernel used local arrays per thread, the compiler will try to store some of them onto registers and spill the ones left to local memory that is also on the DRAM
Stuff Learned till now:
- Global memory coalescing
- Matrix multiplication tiling (indices not yet correctly implemented)
Why dont we do it like every thread in the row loads a single item and just use the blocks, because the shared memry can only be accessed by the same block