Best Practices#

When implementing your simulation, there are some common patterns where each of them have certain best practices that help you to achieve better performance.

Important

PLEASE, PLEASE, PLEASE read this page carefully before doing your implementation. We have seen too many previous benchmark applications with poor performance or race condition due to not following these best practices. And we don’t want to see the same mistakes again.

Prefer Less `callAll` Invocations#

It is true that separating different calculation within a simulation iteration into multiple callAll calls can provide better modularity and readability. However, each callAll invocation introduces one GPU kernel launch, and each kernel has to wait until the previous one fully finishes before it can start spawning threads and executing on GPUs, which can introduce notable latency.

As a result, it is usually more recommended to group as much as you can into a single callAll invocation to minimize such overhead. If you really want to modularize your code, consider implementing private member methods inside your custom Place or Agent class and call them sequentially from a single callAll invocation.

Beware of Race Condition#

MASS CUDA allow each Place instance to access its neighbors directly and allow each Agent instance to access its residing Place instance directly. Although such design can result in less overhead, race condition becomes a concern that you need to be aware of. When implementing your simulation, always remember that each member method you implemented in Place or Agent class will be executed by all the valid instances concurrently and always be careful not to overlap read and write operations on the same attribute.

A common mistake that will lead to race condition is trying to update an attribute base on values of the same attribute of the neighbors. For example, if we have a custom Place class named MyPlace containing an attribute named ATTR, the func method in the following code will lead to race condition:

class MyPlace final : public mass::Place {
public:
    __universal__ explicit MyPlace(const mass::Index localIndex) : Place(localIndex) {}

    enum Attribute { ATTR };
    enum Function { FUNC };

    __device__ void callMethod(const mass::FunctionId functionId, void*const arg) {
        switch (functionId) {
            case FUNC: func(); break;
            default: break;
        }
    }

    __device__ void func() {
        const auto attrPtr = getAttribute<int>(Attribute::ATTR);
        auto newValue = *attrPtr;
        const auto neighborPtrs = getNeighborsPtr();
        for (auto i = 0; i < 4; i++) {
            newValue += *neighborPtrs[i]->getAttribute<int>(Attribute::ATTR);
        }
        *attrPtr = newValue;
    }
}

When running the func method through callAll, all the non-ghost instances run it concurrently. Consider at certain time point, instance $p$ runs into the for-loop trying to access values from its neighbors. At this point, since the execution order is not guaranteed, some of its neighbors may have already finished running the func method and have its ATTR value updated while some other neighbors may not have started yet. As a result, the values gathered by instance $p$ may not be consistent where some of them are updated values and some of them are old values.

The solution to such situation is to use a phase-based double buffering design. We need to declare the ATTR attribute as an 2D attribute of length 2 and use an integer argument ranging between 0 and 1 to determine the current phase. If phase == 0, we use ATTR[0] as the old value and ATTR[1] for the updated value. In contrast, if phase == 1, we use ATTR[1] as the old value and ATTR[0] for the updated value. The simulation loop is responsible for updating the phase value in each iteration. The updated code should look like the following:

// MyPlace.h

class MyPlace final : public mass::Place {
public:
    __universal__ explicit MyPlace(const mass::Index localIndex) : Place(localIndex) {}

    enum Attribute { ATTR };
    enum Function { FUNC };

    __device__ void callMethod(const mass::FunctionId functionId, void*const arg) {
        switch (functionId) {
            case FUNC: func(*static_cast<const int*>(arg)); break;
            default: break;
        }
    }

    __device__ void func(const int phase) {
        const auto attrPtr = getAttribute<int>(Attribute::ATTR);
        auto newValue = attrPtr[phase];
        const auto neighborPtrs = getNeighborsPtr();
        for (auto i = 0; i < 4; i++) {
            newValue += neighborPtrs[i]->getAttribute<int>(Attribute::ATTR)[phase];
        }
        attrPtr[1 - phase] = newValue;
    }
}

// main.cpp

places->setAttribute<int>(
    /* tag = */ MyPlace::Attribute::ATTR,
    /* elementCount = */ 2,
    /* initialValue = */ 1
);
places->finalizeAttributes();

for (auto i = 0; i < iterations; i++) {
    const auto phase = i % 2;
    places->callAll(MyPlace::Function::FUNC, phase);
}

In this new implementation, because we separate the read source and write destination into different indices of the attribute base on the phase value, the race condition is eliminated.

Choose Your Data Storage Method Wisely#

For storing your data during simulation for each instance, choose the method wisely. Do not try to store everything as normal attributes since that may lead to higher bandwidth usage and worse performance. We have a few rules for picking the right method for storing your data:

If your data is shared across all instances and stays constant during the simulation process, use Shared Attributes. Technically you can also use callAll argument for this, but adding too many things to callAll argument can make the code a little ugly while not providing notable performance improvement.
If your data is shared across all instances but may change during the simulation process, use callAll argument. In this case, prepare your data on host-side code and pass it to callAll invocation. If you have multiple values to pass, use a custom struct or a tuple to wrap them up. The phase value in the Beware of Race Condition section is a perfect example of this situation.
If your data can have different values for different instances, use Normal Attributes. Normal Attributes are way slower than the other two, but in such situation, this is the only choice.

Minimize `syncPlaceBorder` Invocation and Include only Necessary Attributes#

The syncPlaceBorder will transfer values of all included attributes from the instances at the boundary of a block to the corresponding Ghost Regions on neighboring devices. This is a time-consuming process, so you should try to minimize the frequency of invoking it and minimize the amount of data being transferred by reducing the number of attributes included.

There are some general guidance for achieving this goal:

Only insert invocation of syncPlaceBorder before calling a callAll that will access attribute values from neighboring Place instances.
Only include attributes that will be accessed from neighboring instances within the subsequent callAll invocation.
Only include attributes that are modified since the previous syncPlaceBorder invocation.

Note that these rules does not guarantee to provide the optimal performance in all simulation, so you should always analyze your implementation to see whether it can be further optimized.

Avoid Branching on Host-Specific Data#

When running simulation on multiple hosts, avoid writing branching (e.g.: if-else, while-loop) depending on host-specific data. The Host-Specific data here refer to data with different values on different hosts, such as rank of host and current time. Doing so may cause processes on different hosts run into different branches and invoke different MASS CUDA APIs and lead to undefined behavior.

The general guidance here is always trying to make sure the branching condition produce the same result across all hosts. If for some reason you have to use host-specific data for branching (e.g.: gathering data of all processes to the master node), make sure not to include any MASS CUDA API calls that involve cross-device communication within each branch.

Configure the Library Based on Your Need#

The MASS CUDA library provides several configurable build options and correct configuration can lower unnecessary memory usage. The core idea is to only include what you need and exclude what you don’t. For example, if your simulation does not use Agents at all (e.g.: Heat2D), set MASS_MAX_AGENTS and MASS_MAX_POTENTIAL_AGENTS to 1 (Currently we don’t support setting them to 0. Should be considered in future release). In addition, set MASS_MAX_NEIGHBORS to exactly the amount you need.

Write Codes as a C++ Developer, Not C#

We are using C++ and CUDA C++ as our source languages for using MASS CUDA. So please try to write codes in C++ standard and style. Use modern C++ features and APIs such as std::vector, std::unique_ptr and references instead of manual memory management and raw pointers.

It is not forbidden to use things like raw pointers, it is still useful in some situation, especially the void*, which is important for doing things like type erasure. However, the key point here is to fully understand what you are doing and only use such unsafe C features when it’s truly necessary.

Best Practices

Contents