The machine: 2xIntel(R) Xeon(R) Gold 6326 CPU @ 2.90GHz, 2xNVIDIA A100 80GB PCIe, no p2p GPU access.

AMD CPU server: 2xAMD EPYC 9754

Periodic shock box

Test solver. 2nd order + Barth limiter + SSPRK3, CFL=0.5

Number of cells iterations per second, CI/s

Using 1024^2 cells.

Using 16 OMP thread x 2 ranks performs nearly the same as (slightly worse than) 32 ranks.

GPU python profile results

CPU python profile results

NSYS results

HUGE RX (host to device memcpy), 6-8GB/s why?

GPU occupancy (nvtop): ~60-70%!

Problem: unintended to_device() calls in initializing (rechecking) face buffer.

Fixing extra to_device

Fixed:

GPU occupancy (nvtop): ~90%, RX/TX several MB/s

2 GPU v.s. 1 GPU: 88% strong scaling efficiency.

Optimized RecGradient and RecFace2nd

Primary optimization: local cache.

When only optimize RecGradient, consider the effect of using shared shuffle write or not:

Shared write (to 3x5 gradient) vs. direct write:

• total: 36.4MCI/s vs. 35.8MCI/s
• RecGradient: 2677 Iter, 8.79s vs/ 9.99s

Around 10% improvement.

Pitfall

If we use a buffer write function using __shared__ + __syncthreads() inside, you might want:

1
2
3
4
5
6
7
8

int tid_global = blockDim.x * blockIdx.x + threadIdx.x; 
if (tid_global >= max) 
{ 
  write_data(dummy_data); 
  return; 
} 
do_calculation.... 
write_data(real_data);

To handle OOB threads. If write_data is templated or inlined, the __shared__ buffer could diverge.

Safe pattern:

1
2
3
4
5
6
7

int tid_global = blockDim.x * blockIdx.x + threadIdx.x; 
t_buffer real_data;
if (tid_global < max) 
{ 
  do_calculation.... 
} 
write_data(real_data);

When both of RecGradient and RecFace2nd are optimized, performance:

Occupancy: 86% 1 GPU / 71% 2 GPU

Write coalescing optimized

Occupancy: 83% 1 GPU / 66% 2 GPU

Larger case 256^3 (3D, more work per cell)

Occupancy: 99% 1 GPU / 90% 2 GPU