Outline:

1. Load balancing intro
2. Dynamic load balancing in
   PIC code run on GPUs

GPU-accelerated machines entered the TOP500 rankings just over a decade ago.

How do we get optimal performance from these supercomputers?

• Compilers
• Algorithms/data structures
• Load balancing

Particle-mesh codes parallelize via domain decomposition.

Particle-mesh codes parallelize via domain decomposition.

Particle-mesh codes parallelize via domain decomposition.

Particle-mesh codes parallelize via domain decomposition.

0

1

1

0

0

1

1

0

Particle-mesh codes parallelize via domain decomposition.

0

1

1

0

Compute
work

Rank

0

1

1

0

Particle-mesh simulations can suffer from load imbalance.

0

1

1

0

Compute
work

Rank

0

1

1

0

Particle-mesh simulations can suffer from load imbalance.

0

1

1

1

Compute
work

Rank

0

1

1

0

Load imbalance can be corrected at run time.

Basic load balance algorithm for distributed memory particle-mesh:

              if (step % loadBalanceInterval == 0) {
                  float currEff = 0.0, propEff = 0.0;
                  DistMapping newDM = makeNewDM(costs,
                                                currEff, propEff);
                  bool globUpdateDM = false;
                  if (myRank == root) {
                      globUpdateDM = (propEff > 1.1*currEff);
                  }
                  bcast(&globUpdateDM, 1, root);
                  if (globUpdateDM) {
                      bcast(&newDM[0], newDM.size(), root);
                      updateDistributionMapping(newDM);
                  }
              }
              

How should costs be measured when running on a GPU-accelerated machine?

1. Start timer
2. Launch kernel
3. Stop timer

How should costs be measured when running on a GPU-accelerated machine?

1. Start timer
2. Launch kernel
3. Stop timer

These are a few strategies appropriate for cost assessment on GPU machines:

• Heuristic: number of particles and cells as proxy for compute work
• CUPTI: use CUDA Profiling Tools Interface to access kernel times
• GPU clock: use thread-summed kernel times as relative measure of compute work

How to measure costs with heuristic?

Cost for rank $i$ is linear combination of number of particles and cells:

\[\begin{aligned} c_{i} = \alpha \cdot n_{\rm particles} + \beta \cdot n_{\rm cells} \end{aligned} \]

• $\alpha$ and $\beta$ are parameters representing relative computational cost of single particle vs. single cell
• $\alpha$ and $\beta$ change depending on algorithm, hardware
• In general, $\alpha$ and $\beta$ should be measured
• Pros: vendor agnostic, low overhead
• Cons: cumbersome tuning of parameters

How to measure costs with CUPTI?

• Pros: API enables powerful profiling capabilities
• Cons: overhead, vendor specific

How to measure costs with CUPTI?

Initialize the trace:

                  cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL);
                  cuptiActivityRegisterCallbacks(bfrRequest, bfrCompleted);
              

Trigger callback functions:

                  void CUPTI API bfrRequest (uint8_t **bfr, ...)
                  {
                      // Signal to CUPTI client that an empty buffer
                      // is needed by CUPTI
                  }
                  void CUPTI API bfrCompleted (uint8_t *bfr, ...)
                  {
                      // Return a buffer of completed activity records
                      // to CUPTI client  
                  }
              

How to measure costs with CUPTI?

Initialize the trace:

                  cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL);
                  cuptiActivityRegisterCallbacks(bfrRequest, bfrCompleted);
              

Trigger callback functions:

                  void CUPTI API bfrRequest (uint8_t **bfr, ...)
                  {...}
                  void CUPTI API bfrCompleted (uint8_t *bfr, ...)
                  {...}
                  
                  ⋮
                  
                  mykernel<<<...>>>(...);
                  cuptiActivityFlushAll(0); // Wait for return of CUPTI 
                  ➞ bfrCompleted(...);     // records via callback function
              

How to measure costs with GPU clock?

• Pros: vendor agnostic, no hyperparameter tuning
• Cons: requires some data movement

How to measure costs with GPU clock?

Add the thread cycles, using atomicAdd for thread safety:

                  __global__ void mykernel (...) {
                      float cycles = clock();
                      ⋮
                      // thread work
                      ⋮
                      cycles = clock() - cycles;

                      // cost_ptr is the pointer to rank's cost
                      atomicAdd(cost_ptr, cycles);
                  }
              

• Reduced overhead using pinned host memory
• To use this: instrument most expensive kernels
• Overcomes weakness of heuristic: that has no sensitivity to how much particles move

Outline:

1. Load balancing intro
2. Dynamic load balancing in
   PIC code run on GPUs

We studied these strategies in the particle-in-cell code WarpX.

We choose laser-ion acceleration as a challenging test problem.

The inhomogeneity translates to different computational costs.

Computational costs are used to compute optimal mapping from MPI rank to domain.

Dynamic load balancing is crucial to performance.

With optimal selection of parameters, we achieve around 3x–4x speedup.

How much improvement expected from load balancing?

The load balancing scheme achieves 62%–74% of theoretical maximum.

What is new/innovative about this work?

With new strategies for GPU cost assessment, we achieved 3x–4x speedup on challenging plasma physics problem.

Thank you! I am happy to answer any questions.

Performance is tuned with additional algorithm-specific parameters.