Repository: ggml-org/llama.cpp
Pull request: Fix data race in CUDA's "cpy" kernel (influences GGML's DUP, CONT operations).

# What is it?

There is a data race at particular kernel, which influences DUP and CONT operations. 
See [ggml/src/ggml-cuda/cpy.cu](https://github.com/Exile333/llama.cpp/blob/master/ggml/src/ggml-cuda/cpy.cu#L61-L88) at lines 61-88. In this loop, one first loads a part of input matrix into shared memory, then loads the shared memory's contents into output matrix. There is a barrier synchronization after the data is loaded. However, it is not enough to prevent data race.

Let's look at the following situation. Both input and output matrices' shape is [10, 10, 5, 1]. This kernel will iterate threadblock over 5 matrices of shape [10, 10]. In the current master branch, one may encounter situation when some warp have read matrix #i, waited all other warps at the barrier, written data to the output matrix #i and went on to reading the next matrix #{i+1}. Given the non-trivial threads' assignment to in/out matrices' coordinates, this will lead to situation where one warp have read portion of the next matrix and another warp writes this data to the previous matrix.

The solution is to add another barrier after the data store operation.

# Ways to test and reproduce.

I've caught this error on AMD Radeon PRO W7900 GPU, in Release build. In all other builds error, caused by data race, occurs much less frequently.

```bash
git clone https://github.com/ggml-org/llama.cpp.git
HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" cmake -B build -DGGML_HIP=ON -DGGML_HIP_ROCWMMA_FATTN=OFF -DGGML_OPENMP=OFF -DAMDGPU_TARGETS=gfx1100 -DCMAKE_C_COMPILER=amdclang -DCMAKE_CXX_COMPILER=amdclang++ -DCMAKE_BUILD_TYPE=Release && cmake --build build --config Release -- -j 16
./build/bin/test-backend-ops test -o "DUP(type=f16,ne=[10,10,5,1],permute=[1,0,2,3])"
```

*NOTE*: as this error is caused by data race, you may need a several launches before it happens.

# Additional tests.

Not needed, current ones are good.

Changed files:
filename: ggml/src/ggml-cuda/cpy.cu
patch: @@ -56,7 +56,8 @@ static __global__ void cpy_scalar_transpose(const char * cx, char * cdst, const
     const int tx = blockIdx.y * CUDA_CPY_TILE_DIM_2D + threadIdx.x;  // transpose block offset
     const int ty = blockIdx.x * CUDA_CPY_TILE_DIM_2D + threadIdx.y;
 
-    __shared__ float tile[CUDA_CPY_TILE_DIM_2D][CUDA_CPY_TILE_DIM_2D+1];
+    __shared__ float tile[2][CUDA_CPY_TILE_DIM_2D][CUDA_CPY_TILE_DIM_2D+1];
+    int cur_tile_buf = 0;
 
 #pragma unroll
     for (int i = 0; i < CUDA_CPY_BLOCK_NM; ++i) {
@@ -70,7 +71,7 @@ static __global__ void cpy_scalar_transpose(const char * cx, char * cdst, const
             if(x < ne01 && y + j < ne00){
                 const int row = threadIdx.y+j;
                 const int col = threadIdx.x * sizeof(float)/sizeof(T);
-                T *tile2 = reinterpret_cast<T*>(tile[row]);
+                T *tile2 = reinterpret_cast<T*>(tile[cur_tile_buf][row]);
                 tile2[col] = src[imat*n + (y+j)*ne01 + x];
             }
         }
@@ -81,10 +82,12 @@ static __global__ void cpy_scalar_transpose(const char * cx, char * cdst, const
         for (int j = 0; j < CUDA_CPY_TILE_DIM_2D; j += CUDA_CPY_BLOCK_ROWS) {
             if (ty + j < ne01 && tx < ne00) {
                 const int col = (threadIdx.y+j)*sizeof(float)/sizeof(T);
-                const T *tile2 = reinterpret_cast<const T*>(tile[threadIdx.x]);
+                const T *tile2 = reinterpret_cast<const T*>(tile[cur_tile_buf][threadIdx.x]);
                 dst[imat*n + (ty+j)*ne00 + tx] = tile2[col];
             }
         }
+
+        cur_tile_buf = (cur_tile_buf + 1) % 2;
     }
 
     GGML_UNUSED_VARS(ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11,