GPU Architecture · Code example

Reverses each block's segment through shared memory, which only works because of __syncthreads(); a -DSKIP_SYNC build plus compute-sanitizer racecheck expose the race when the barrier is removed.

// 02-syncthreads-reverse.cu — see why __syncthreads() is required.
//
// Each block loads its segment into shared memory, then writes it back reversed.
// The reversed write reads a *neighbor's* shared slot, so it must not run until
// the whole tile is loaded — that's the barrier's job.
//
//   Build (correct):        nvcc -arch=sm_120 02-syncthreads-reverse.cu -o rev
//   Build (barrier removed): nvcc -arch=sm_120 -DSKIP_SYNC 02-syncthreads-reverse.cu -o rev_bad
//   Run:                    ./rev      (expect: PASS)
//                           ./rev_bad  (may PASS by luck or FAIL — a data race is UB)
//
// Tip: run either under  compute-sanitizer --tool synccheck ./rev   (barrier checks)
//                   or   compute-sanitizer --tool racecheck ./rev_bad (shared races)

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CHECK(call) do {                                              \
    cudaError_t _e = (call);                                          \
    if (_e != cudaSuccess) {                                          \
      printf("CUDA error at line %d: %s\n", __LINE__,                 \
             cudaGetErrorString(_e));                                 \
      return 1;                                                       \
    }                                                                 \
  } while (0)

#define BLOCK 256

// Reverse each block's own segment of `n` elements via shared memory.
__global__ void reverseBlock(float* out, const float* in, int n) {
  __shared__ float tile[BLOCK];
  int t          = threadIdx.x;
  int blockStart = blockIdx.x * blockDim.x;
  int gid        = blockStart + t;

  if (gid < n) tile[t] = in[gid];          // each thread writes ITS slot

#ifndef SKIP_SYNC
  __syncthreads();                         // wait until the whole tile is written…
#endif

  int rev = blockDim.x - 1 - t;            // …because now we read a NEIGHBOR's slot
  if (gid < n) out[gid] = tile[rev];
}

int main() {
  const int n = 64 * BLOCK;                // exact multiple of BLOCK → no partial blocks
  size_t bytes = (size_t)n * sizeof(float);

  float *hin = (float*)malloc(bytes), *hout = (float*)malloc(bytes);
  for (int i = 0; i < n; ++i) hin[i] = (float)i;

  float *din, *dout;
  CHECK(cudaMalloc(&din, bytes));
  CHECK(cudaMalloc(&dout, bytes));
  CHECK(cudaMemcpy(din, hin, bytes, cudaMemcpyHostToDevice));

  reverseBlock<<<n / BLOCK, BLOCK>>>(dout, din, n);
  CHECK(cudaGetLastError());
  CHECK(cudaDeviceSynchronize());

  CHECK(cudaMemcpy(hout, dout, bytes, cudaMemcpyDeviceToHost));

  // Verify: within each block, out[blockStart + t] == in[blockStart + (BLOCK-1-t)].
  int bad = 0;
  for (int b = 0; b < n / BLOCK; ++b)
    for (int t = 0; t < BLOCK; ++t) {
      float expect = hin[b * BLOCK + (BLOCK - 1 - t)];
      if (hout[b * BLOCK + t] != expect) { bad++; }
    }
  printf("%s  (%d mismatches out of %d)\n", bad == 0 ? "PASS" : "FAIL", bad, n);

  free(hin); free(hout);
  cudaFree(din); cudaFree(dout);
  return 0;
}