Global memory is the GPU's main DRAM. It is what cudaMalloc returns, it is visible to every thread in every block, and it persists between kernel launches. It is also the slowest memory on the device, with latency in the range of 400 to 800 cycles.

float *d_a;
cudaMalloc(&d_a, n * sizeof(float));
cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
kernel<<<blocks, threads>>>(d_a);
cudaFree(d_a);

The headline number on a GPU spec sheet is global memory bandwidth, not FLOPS, and for good reason. Most real kernels are bandwidth-bound, so the achievable fraction of that bandwidth sets the performance ceiling. Measuring it is straightforward: count the bytes the kernel must read and write, divide by the kernel time from CUDA events, and compare against the spec.

Latency is high but throughput is enormous, and the hardware is built to trade one for the other. While one warp waits on a load, the multiprocessor issues instructions from other resident warps. Enough warps in flight hides the latency completely, which is why occupancy matters and why GPUs need far more parallelism than a CPU to reach peak.

What does not get hidden is wasted bandwidth. Memory is served in 32-byte sectors, so a warp reading scattered addresses pulls in far more data than it uses:

a[i]           // stride 1: neighbouring threads hit the same sectors
a[i * 16]      // stride 16: each thread touches its own sector, 16x the traffic

Getting the access pattern right is coalescing, and it is normally the first thing to fix in a slow kernel. The second is to stop going to global memory at all, by staging reused data in shared memory.