Lesson 2: Figuring out LLNL Compilers

This module is designed to help you understand which modules, compilers, and flags you need to use when manually compiling programs on LLNL systems. While most of the time you will be using a build-tool (make, CMake, Spack, etc) that will take care of determining these for you, doing these examples will help you understand the command error messages that pop up when your don't have the correct modules loaded.

Tip

Before even starting to compile, it is always a good idea to double check what modules you have loaded, if not fully reset your session if something seems like it should be working, but is not. Many errors can be caused by the wrong module being loaded (or unloaded), especially when mixing in non-default versions of modules.

The first handful of steps will provide a bare-bones program to compile and run, focusing on figuring out what modules, libraries, and flags are needed to compile a file. The later examples bring in CMake, which has multiple ways to do things, especially for compiling files with GPU code. These steps also should not require a specific compiler or MPI implementation, so feel free to try out multiple compilers and MPI implementations as well. For those testing on LLNL's Dane, only Steps 0 and 1 are possible, as Dane is not a GPU system. Although the GPU examples are geared towards HIP and RCCL, they should, in theory, also work on NVIDIA GPUs (though some may require name changes to the functions used).

Tip

It may seem complicated at first, especially for GPU programs, but LLNL does provide ample documentation on the available compilers in their Livermore Computing PSAAP IV Quick Start Guide.
Furthermore, it is a very good idea to pay attention to the versioning use in both this Module and any documentation found online, as what may have worked in the past may NOT work in the future, or may have a simpler option available.

Step 0: Figuring out the Compilers

The first part of compiling code on an HPC system is figuring out which compiler you want to use. There are often many choices, and these choices will vary from system to system depending on the hardware and system vendors. Often the true compiler will likely be some variant of GCC or Clang, but the wrappers on top of these compilers will still influence how things are compiled.

For example, here is a table of compilers on LLNL's Tuolumne (up to date as of August 10th, 2026), which modules load them, and what compiler they call underneath. Since the point of this lesson is learning about which compilers can automatically handle compiling GPU device code and/or linking to certain libraries, those fields have been hidden with spoiler boxes.

C++ Compiler	Underlying Compiler Family	Module	Compiles GPU Code	Links to MPI	Links to HIP
g++	GNU	(System Default) ****	❌	❌	❌
CC	Cray Clang++	CCE/20.0.0	⚠️*	✅**	⚠️*
mpic++ (mpicxx)	Cray Clang++	cray-mpich/9.0.1	⚠️*	✅	⚠️*
amdclang++	LLVM Clang++	rocm/6.4.3	✅*	❌	✅
hipcc	LLVM Clang++	rocm/6.4.3 or rocmcc/6.4.3	✅	❌	✅
mpiamdclang++	amdclang++	rocmcc/6.4.3 ***	✅*	✅	✅

Show Table Footnotes:

(*) As long as the file is either a `.hip` or `-x hip` is used AND `--offload-arch=gfx942` is used. Without the latter, the code may compile, but it will not successfully run GPU kernels or code.
(**) As long as a `cray-mpich` module is loaded.
(***) Technically this is a wrapper script, and has no version. While you only need to load the `rocmcc` module to see this, it is actually from a specific cray-mpich wrapper package that gets automatically loaded.
(****) Using the GNU-based compilers (including the other versions in the module files) will be "compiling in extreme mode". If the code is only a C/C++ code, then the program _should_ generally work without issues. GPU codes must still be compiled with another (clang-based) option. While it may be technically possible to use a GNU-based compiler for your build process, it will require compiler expertise that is not covered in this module. Direct use of this compiler family on Tuolumne is highly discouraged.

Note

While some compilers may not automatically link to MPI or HIP, it is likely still possible that you can manually do it through the combination of -L, -I, and -l flags. However, if you need to compile GPU device code, your options are more limited, as GPU code has different syntax that a pure C++ compiler may not know how to handle.

Before you get started compiling all of these programs, go ahead and make a directory to pull all of them into. This directory will be helpful for later steps. Once you have that directory made, go ahead and copy, compile, and run the following program:

Click for Step 0 Program

#include<iostream>

int main(int argc, char** argv)
{
    std::cout << "Hello World!" << std::endl;
    return 0;
}

Step 1: Adding in MPI

Compile and run the following MPI program.

Note

Do not run this program on the head node. Most clusters block MPI applications from running on the head node by default. You should run this application on a compute node to verify that you can run it.

Click for Step 1 Program

#include<iostream>

#include "mpi.h"

int main(int argc, char** argv)
{
    MPI_Init(&argc, &argv);

    int rank, size;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    MPI_Comm_size(MPI_COMM_WORLD, &size);

    std::cout << "Hello World from rank " << rank;
    std::cout << " of: " << size << std::endl;

    MPI_Finalize();
    return 0;
}

Step 2: Basic GPU Program

Compile and run the following HIP (ROCm) program. This program does not need any MPI to run, and should run with any number of GPUs.

Tip

The file extension (.cpp, .hip) you use for this program will influence exactly how you need to compile this program. Double check the table in Step 0 if you get stuck figuring out how to compile this code.

When running the program, double check the output you get. You should get five lines of output from the program. If you ended up with only a single line (and got this error: HIP Error: invalid device function at step2.hip:25), then there is likely an additional step you need for compiling GPU-based codes. Check the Table in Step 0 for the compilation flag you are missing.

Additionally, you may want to consider doing module load craype-accel-amd-gfx942 to set some environment variables that may help various toolchains.

Click for Step 2 Program

#include <iostream>
#include <cstdlib>
#include <hip/hip_runtime.h>

#define HIP_CHECK(command) \
    do { \
        hipError_t status = (command); \
        if (status != hipSuccess) { \
            std::cerr << "HIP Error: " << hipGetErrorString(status) \
                      << " at " << __FILE__ << ":" << __LINE__ << std::endl; \
            std::exit(EXIT_FAILURE); \
        } \
    } while (0)

__global__ void helloWorldKernel()
{
    printf("Hello World from GPU thread %d!\n", threadIdx.x);
}

int main(int argc, char** argv)
{
    std::cout << "Hello World from CPU!" << std::endl;

    helloWorldKernel<<<1, 4>>>();
    HIP_CHECK(hipGetLastError());
    HIP_CHECK(hipDeviceSynchronize());

    return 0;
}

Step 3: Mixing MPI and GPU Programming

The next example requires linking with both an MPI library and HIP, which can introduce more complexity in making sure everything is properly linked.

Step 3a: CPU-only MPI Linked with GPU Libraries

Initially, the program will only test to see if you have linked both MPI and HIP. Running this version does not need a GPU-aware MPI implementation.

Click for Step 3 Program

#include <iostream>
#include <vector>
#include <hip/hip_runtime.h>
#include <mpi.h>

// HIP Error checking macro
#define HIP_CHECK(command) \
    do { \
        hipError_t status = (command); \
        if (status != hipSuccess) { \
            std::cerr << "HIP Error: " << hipGetErrorString(status) \
                      << " at " << __FILE__ << ":" << __LINE__ << std::endl; \
            MPI_Abort(MPI_COMM_WORLD, status); \
        } \
    } while (0)

// MPI Error checking macro
#define MPI_CHECK(command) \
    do { \
        int status = (command); \
        if (status != MPI_SUCCESS) { \
            char error_string[MPI_MAX_ERROR_STRING]; \
            int length_of_error_string; \
            MPI_Error_string(status, error_string, &length_of_error_string); \
            std::cerr << "MPI Error: " << error_string \
                      << " at " << __FILE__ << ":" << __LINE__ << std::endl; \
            MPI_Abort(MPI_COMM_WORLD, status); \
        } \
    } while (0)

// Kernel to initialize an array of data
__global__ void initDataKernel(float* data, int rank, int num_elements) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx < num_elements) {
        data[idx] = 1.0f; 
    }

    // Print greeting only from the very first thread to avoid console spam
    if (idx == 0) {
        printf("Hello World from GPU thread 0 on MPI rank %d!\n", rank);
    }
}

int main(int argc, char** argv)
{
    MPI_CHECK(MPI_Init(&argc, &argv));

    int rank, size;
    MPI_CHECK(MPI_Comm_rank(MPI_COMM_WORLD, &rank));
    MPI_CHECK(MPI_Comm_size(MPI_COMM_WORLD, &size));

    const int buffer_size = 1000;
    const size_t bytes = buffer_size * sizeof(float);

    float *d_sendbuf, *d_recvbuf;
    HIP_CHECK(hipMalloc(&d_sendbuf, bytes));
    HIP_CHECK(hipMalloc(&d_recvbuf, bytes));

#ifdef TEST_GPU_AWARE
    int threadsPerBlock = 256;
    int blocks = (buffer_size + threadsPerBlock - 1) / threadsPerBlock;

    initDataKernel<<<blocks, threadsPerBlock>>>(d_sendbuf, rank, buffer_size);

    // Catch any launch errors
    HIP_CHECK(hipGetLastError());
    HIP_CHECK(hipDeviceSynchronize());

    // Execute GPU-aware MPI collective on all 1000 items
    MPI_CHECK(MPI_Allreduce(d_sendbuf, d_recvbuf, buffer_size, MPI_FLOAT, MPI_SUM, MPI_COMM_WORLD));

    // Allocate host memory using std::vector and fetch results
    std::vector<float> h_recvbuf(buffer_size);
    HIP_CHECK(hipMemcpy(h_recvbuf.data(), d_recvbuf, bytes, hipMemcpyDeviceToHost));

    // Verification
    bool correct = true;
    float expected_val = static_cast<float>(size);
    for (int i = 0; i < buffer_size; ++i) {
        if (h_recvbuf[i] != expected_val) {
            correct = false;
            break;
        }
    }

    if (correct) {
        std::cout << "Rank " << rank << ": Verification passed! All " 
                  << buffer_size << " items matched expected value (" << expected_val << ")." << std::endl;
    } else {
        std::cerr << "Rank " << rank << ": Verification FAILED! Data mismatch encountered." << std::endl;
    }
#else
    std::cout << "Hello MPI with HIP from rank " << rank << " of: " << size << std::endl;
#endif

    HIP_CHECK(hipFree(d_sendbuf));
    HIP_CHECK(hipFree(d_recvbuf));
    MPI_CHECK(MPI_Finalize());

    return 0;
}

Step 3b: GPU-Aware MPI Example

Now try re-compiling the above example with -DTEST_GPU_AWARE to your compilation flag. Adding this flag will cause the test to attempt to use a GPU buffer in an MPI communication call. This will produce errors when running if you have not compiled the program correctly.

This type of program is the version most likely to produce problems. While the compilation errors (i.e., <mpi.h> not found, <hip/hip_runtime.h> not found) are fairly simple to solve (using the methods you used in Steps 1 and 2), this variant can cause two types of runtime errors. The first is caused by the lack of GPU-aware support in the MPI implementation. On Tuolumne, the cray-mpich library is GPU-aware, but you will need to set the environment variable MPICH_GPU_SUPPORT_ENABLED before running the program. Without this set, if you attempt to use GPU Aware cray-mpich, it should warn about it and exit. Occasionally it may produce other errors, such as "flux-job: task(s) Bus error".

The second is technically caused by compilation, but the issue does not present itself until you actually run the program. Specifically, you may see this error: "MPIDI_CRAY_init: GPU_SUPPORT_ENABLED is requested, but GTL library is not linked". As the message implies, there is an additional library that cray-mpich specifically needs to link to in order to enable GPU-aware support. A good way to check this before running is to use ldd (explained more in Step 6) and looking for a specific GTL line:

# Example output if library is present:
ldd my_library_or_executable | grep "gtl"
libmpi_gtl_hsa.so.0 => /opt/cray/pe/lib64/libmpi_gtl_hsa.so.0 (0x000015555507a000)
# Example output if library is not present:
ldd my_library_or_executable | grep "gtl"

Adding in the extra library may automatically happen based on the compiler. The following table shows how the compilers found in Step 0 handle this extra library:

Compiler	Automatically Adds GTL library?
CC	⚠️*
mpic++	⚠️**
amdclang++	❌
hipcc	❌
mpiamdclang++	✅

(*) By default, CC will not link with the GTL library. However, if you add --offload-arch=gfx942 when compiling, or do module load craype-accel-amd-gfx942 before compiling, the CC compile will automatically grab the GTL library. Given that one or both are best practice (and explained above), the box is marked "yes".

(**) By default, the cray-provided mpic++ command does not automatically link with the GTL library. However, the Tuolumne system administrators have added custom mpic++ wrappers that automatically add the GTL library to the compilation. These Tuo-specific scripts are loaded when you load any rocmcc module (but not with the rocm module). For the default mpic++, you simply only need to add the flag -lmpi_gtl_hsa to link with the correct GTL library (CUDA systems should use -lmpi_gtl_cuda).

Note

Both of the above issues are specific to cray-mpich. Other MPI implementations may have different toggles for GPU-aware support, so consult the documentation before attempting to use the above strategies on different MPI implementations.

If you must use either hipcc or amdclang++, the -lmpi_gtl_hsa is still the correct addition. However, you will still need the appropriate -L, -I, -l for the regular MPI library. The -L for the MPI library will serve double duty for the -lmpi_gtl_hsa as well since cray puts both in the same directory.

Step 4: Using RCCL

RCCL (or ROCm Communication Collectives Library), is the ROCm equivalent for NCCL (NVIDIA Communication Collectives Library). Both of these are similar to MPI, but designed for communication between GPUs. This Module will not focus on understanding how these libraries work or how to build them. Instead it focuses on compiling and running them.

To help you get started with compilation, you will need to add -lrccl, as by default, the RCCL library is not automatically linked to by any compiler. The header (-I) and library (-L) paths should not need to be added, but the compiler will tell you if they need to be with something like "rccl/rccl.h cannot be found" or "librccl.so cannot be found"

Note

While HIP is GPU agnostic, all of the RCCL calls will have the same "nccl" prefix as if you were using NCCL.

There is also not a "RCCL" specific module on Tuolumne that you need to load. RCCL gets loaded with the ROCm module. The RCCL modules that are present on Tuolumne are for tweaking how RCCL runs, which will not be covered here.

Tip

This test is very sensitive to the values of HIP_VISIBLE_DEVICES (or ROCR_VISIBLE_DEVICES). When choosing how to run this test, you may want to run this test on an interactive session without flux run (but don't run on the login node!). From our testing, you will need to use --exclusive on your Flux allocation in order to make sure flux gives enough GPUs.

Click for Step 4 Program

#include <iostream>
#include <vector>
#include <hip/hip_runtime.h>
#include <rccl/rccl.h>

int main() {
    int numGpus = 0;
    hipGetDeviceCount(&numGpus);

    if (numGpus < 2) {
        std::cout << "Error: This example requires at least 2 AMD GPUs." << std::endl;
        return 1;
    }

    const int nDev = 2;
    int devs[nDev] = {0, 1};

    ncclComm_t comms[nDev];
    hipStream_t streams[nDev];
    float* d_data[nDev];

    // Data to sum across GPUs: GPU 0 sends 10.0, GPU 1 sends 20.0
    float h_data[nDev] = {10.0f, 20.0f};
    float h_result[nDev] = {0.0f, 0.0f};

    // 1. Initialize RCCL communicators for the 2 GPUs
    ncclCommInitAll(comms, nDev, devs);

    // 2. Allocate GPU memory & set up streams
    for (int i = 0; i < nDev; ++i) {
        hipSetDevice(devs[i]);
        hipStreamCreate(&streams[i]);
        hipMalloc(&d_data[i], sizeof(float));
        hipMemcpyAsync(d_data[i], &h_data[i], sizeof(float), hipMemcpyHostToDevice, streams[i]);
    }

    // 3. Launch the RCCL collective operation (Sum AllReduce)
    // Group calls let RCCL coordinate execution across devices cleanly
    ncclGroupStart();
    for (int i = 0; i < nDev; ++i) {
        hipSetDevice(devs[i]);
        ncclAllReduce((const void*)d_data[i], (void*)d_data[i], 1,
                      ncclFloat, ncclSum, comms[i], streams[i]);
    }
    ncclGroupEnd();

    // 4. Synchronize streams & copy results back to CPU
    for (int i = 0; i < nDev; ++i) {
        hipSetDevice(devs[i]);
        hipStreamSynchronize(streams[i]);
        hipMemcpy(&h_result[i], d_data[i], sizeof(float), hipMemcpyDeviceToHost);

        std::cout << "GPU " << devs[i] << " received AllReduce result: " 
                  << h_result[i] << std::endl;
    }

    // 5. Clean up resources
    for (int i = 0; i < nDev; ++i) {
        hipFree(d_data[i]);
        hipStreamDestroy(streams[i]);
        ncclCommDestroy(comms[i]);
    }

    return 0;
}

Step 5: Using RCCL and MPI

This step breaks up compilation and running, as the latter is more involved than the previous setups, and may introduce new errors.

Step 5a: Compilation

Compiling this program requires combining the methods used in the two previous steps, but otherwise should not add anything new. Note that this example should not produce any warnings, as wrappers for checking the return codes have been introduced. These will be helpful for debugging any issues that pop up while running.

Tip

Start with the command you used in Step 3, and add the flag(s) you used in Step 4.

Click for Step 5 Program

#include <iostream>
#include <cstdlib>
#include <mpi.h>
#include <hip/hip_runtime.h>
#include <rccl/rccl.h>

#define HIPCHECK(cmd) do { \
    hipError_t e = cmd; \
    if (e != hipSuccess) { \
        std::cerr << "HIP Error: " << hipGetErrorString(e) << " at line " << __LINE__ << std::endl; \
        MPI_Abort(MPI_COMM_WORLD, e); \
    } \
} while(0)

#define RCCLCHECK(cmd) do { \
    ncclResult_t r = cmd; \
    if (r != ncclSuccess) { \
        std::cerr << "RCCL Error: " << ncclGetErrorString(r) << " at line " << __LINE__ << std::endl; \
        MPI_Abort(MPI_COMM_WORLD, r); \
    } \
} while(0)

int main(int argc, char* argv[]) {
    MPI_Init(&argc, &argv);

    int rank, size;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    MPI_Comm_size(MPI_COMM_WORLD, &size);

    // 1. Check what GPUs the job scheduler is exposing to this specific rank
    const char* hip_vis = std::getenv("HIP_VISIBLE_DEVICES");
    const char* rocr_vis = std::getenv("ROCR_VISIBLE_DEVICES");

    std::cout << "Rank " << rank << " sees HIP_VISIBLE_DEVICES=" 
              << (hip_vis ? hip_vis : "unset") 
              << " | ROCR_VISIBLE_DEVICES=" 
              << (rocr_vis ? rocr_vis : "unset") << std::endl;

    // Wait for all ranks to print their env variables before proceeding
    std::flush(std::cout);
    MPI_Barrier(MPI_COMM_WORLD);

    // 2. Calculate LOCAL rank on this specific node
    MPI_Comm localComm;
    MPI_Comm_split_type(MPI_COMM_WORLD, MPI_COMM_TYPE_SHARED, rank, MPI_INFO_NULL, &localComm);

    int localRank;
    MPI_Comm_rank(localComm, &localRank);

    // 3. Map local rank to a GPU
    int numGpus = 0;
    HIPCHECK(hipGetDeviceCount(&numGpus));

    if (numGpus == 0) {
        std::cerr << "Rank " << rank << ": No GPUs found on node." << std::endl;
        MPI_Abort(MPI_COMM_WORLD, 1);
    }

    // Determine GPU mapping. If the scheduler is isolating GPUs (numGpus == 1), 
    // we must use device 0. Otherwise, we distribute based on localRank.
    int localDev = (numGpus == 1) ? 0 : (localRank % numGpus);
    HIPCHECK(hipSetDevice(localDev));

    std::cout << "Rank " << rank << " bound to local device ID " << localDev 
              << " (Total visible to this rank: " << numGpus << ")" << std::endl;
    std::flush(std::cout);
    MPI_Barrier(MPI_COMM_WORLD);

    // 4. Broadcast RCCL Unique ID from Rank 0
    ncclUniqueId id;
    if (rank == 0) {
        RCCLCHECK(ncclGetUniqueId(&id));
    }
    MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);

    // 5. Initialize RCCL Communicator
    ncclComm_t comm;
    RCCLCHECK(ncclCommInitRank(&comm, size, id, rank));

    // 6. Create Stream & Memory
    hipStream_t stream;
    HIPCHECK(hipStreamCreate(&stream));

    float* d_data = nullptr;
    HIPCHECK(hipMalloc(&d_data, sizeof(float)));

    float h_data = (rank + 1) * 10.0f;
    float h_result = 0.0f;

    HIPCHECK(hipMemcpyAsync(d_data, &h_data, sizeof(float), hipMemcpyHostToDevice, stream));

    // 7. Launch AllReduce
    RCCLCHECK(ncclAllReduce((const void*)d_data, (void*)d_data, 1, 
                            ncclFloat, ncclSum, comm, stream));

    // 8. Synchronize and Copy back
    HIPCHECK(hipStreamSynchronize(stream));
    HIPCHECK(hipMemcpy(&h_result, d_data, sizeof(float), hipMemcpyDeviceToHost));

    // Print final result
    std::cout << "Rank " << rank << " received AllReduce result: " << h_result << std::endl;

    // Clean up
    HIPCHECK(hipFree(d_data));
    HIPCHECK(hipStreamDestroy(stream));
    RCCLCHECK(ncclCommDestroy(comm));
    MPI_Comm_free(&localComm);
    MPI_Finalize();

    return 0;
}

Step 5b: Running

In order to run this application, you will need to make sure two things happen: Either HIP_VISIBLE_DEVICES or ROCR_VISIBLE_DEVICES is set, and no two MPI ranks share the same device. A given MPI rank may have multiple devices, but they cannot share devices between ranks. This may require you to look into the flux run command documentation in order to get very specific on how the GPUs are assigned. The program above will print out what it sees so you can check. The final result of the program will depend on how many ranks were involved. One rank should produce a value of 10, two ranks should produce a value of 30, three ranks a value of 60, and four rank a value of 100.

You do not have to run this program on multiple nodes, but it should still work if you choose to do so.

Tip

Getting an interactive allocation is still a good idea, as you can play around with many different flux run attempts in a single go. It is also a good idea to add --exclusive to your allocation (either flux alloc or flux batch) to ensure that Flux does not attempt to divide up the node you are given.
Additionally, there is more nuance on how Flux will allocate GPUs to a given run. Using the -n flag will always result in 1 GPU for the entire run, regardless of the value given to -n and the value given to -g. Instead, we currently recommend using --tasks-per-node and -N on a flux run to easily get the number of desired GPUs.

Use of the flag -o mpibind=verbose:1 on flux run command is highly recommended when attempting to manage GPU mappings for the first time.

Step 6: Checking the Libraries and Mixing Modules

While most programs will not need to load a non-default GPU library, there is still an occasional need to load a different one. This may be because cutting edge research needs new features, bug fixes, or other compatibility reasons. Regardless, loading modules introduces more avenues for the "wrong library version" headache. This section focuses on how you can look out for these issues to try and catch them earlier.

The first step when mixing module/library version is to interrogate the compiler itself. This can be done us commands, such as which, to double check that the compiler has changed. While often times a compiler will change when you load a different module for it, it may stay the same if it is actually a generic wrapper (which does happen LLNL systems). This is when asking the various build tools to output their versions, or increase their verbosity, can be helpful. For example --version on a tool is always a quick and simple sanity check; make VERBOSE=1 is a great way to see which libraries are being pulled in; And some compilers even offer a --verbose flag to make the compilation itself more verbose in what it is doing.

However, just because all of the tools and output above looks correct, there may still be hidden issues that can arise. One additional tool you can use is ldd. This tool will show you which shared libraries your (shared) library or executable will attempt to load. If any of these look off, then it means either you don't have the right modules loaded when you ran ldd, or the compilation itself was not correct. Use ldd on any of your executables from above and take a look at what libraries are linked into your program.

Note

ldd only works on shared libraries (.so) or programs built in a non-static way. A static (.a) library will need to use different approach, or watch the output of the compilation very carefully.

Additionally, is it possible for the compiler to hard-code a .so path into your program, which can make it more resilient to module switches. Exactly how to do that will not be covered here.

For reference, here is what ldd looks like when you compile with a module loaded, but then don't have it loaded for ldd:

    libmpi_gtl_hsa.so.0 => /opt/cray/pe/lib64/libmpi_gtl_hsa.so.0 (0x000015555507a000)
    libmpi_cray.so.12 => /opt/cray/pe/lib64/libmpi_cray.so.12 (0x0000155553f7e000)
    libamdhip64.so.6 => not found
    libcrayacc_amdgpu.so.1 => /opt/cray/pe/cce/20.0.0/cce/x86_64/lib/libcrayacc_amdgpu.so.1 (0x00001555527d8000)
    libamdhip64.so.6 => not found
    libhsa-runtime64.so.1 => not found
    libamdhip64.so.6 => not found

If there are every any "not found" when your program tries to run, it will crash because it could not load the libraries it needs to do. That's why using ldd with the module setup you expect to run with is a good way to test things before submitting a flux job.

And then when the correct modules are loaded when ldd tries to run:

    libmpi_gtl_hsa.so.0 => /opt/cray/pe/lib64/libmpi_gtl_hsa.so.0 (0x000015555507a000)
    libmpi_cray.so.12 => /opt/cray/pe/lib64/libmpi_cray.so.12 (0x00001554d35c7000)
    libamdhip64.so.6 => /opt/rocm-6.4.3/lib/libamdhip64.so.6 (0x00001554d0633000)
    libcrayacc_amdgpu.so.1 => /opt/cray/pe/cce/20.0.0/cce/x86_64/lib/libcrayacc_amdgpu.so.1 (0x00001554d04d7000)
    libhsa-runtime64.so.1 => /opt/rocm-6.4.3/lib/libhsa-runtime64.so.1 (0x00001554cfbc9000)
    libamd_comgr.so.3 => /opt/rocm-6.4.3/lib/libamd_comgr.so.3 (0x00001554c4f46000)
    libdrm_amdgpu.so.1 => /lib64/libdrm_amdgpu.so.1 (0x00001554c4700000)

As a final parting thought, consider using grep to help search the output of ldd, using something like this:

ldd <my_executable-or_library> | grep -E "(amd|gtl|mpi|hsa)"

This command will search for any line that has "amd", "gtl", "mpi", or "hsa" (underlying ROCm library). Remember that you do not want to mix multiple GPU libraries, so if the same library shows up twice (such as having both libamdhip64.so.6 and libamdhip64.so.7), or some GPU libraries have one path while others have a different path, then something has gone wrong in the compilation step.

For this step, play around with each of the above commands and flags to see what type of output you get (and save make VERBOSE=1 for Step 7). You do not have to load a non-default module to complete this step, but doing so may help reveal what issues arise when you do so.

Step 7: Adding in CMake

All of the previous steps should cover most of the basic use cases for compiling MPI and GPU programs. There may be other GPU libraries that need more specific compilation flags, but the process should align to one of the previous examples. This section will take that understanding to help debug CMake when things go wrong.

Warning

While this section does not attempt to introduce more complex CMake than is necessary, it does assume a basic knowledge of how CMake roughly works. If you need a refresher, consider taking a look at our CMake Articles!

Note

This step, and any step that follows, will be much more tailored to AMD GPUs than before. While it should still be closer to a one-to-one conversion, it will be more involved than simply tweaking the programs.

Step 7a: Setting up CMake

At this point you should have multiple files in your current directory. You will now be adding one more file, but this time it will be a CMake file below. Add this file to the directory, call it CMakeLists.txt, and then fill in any <ADD ME!> sections with the appropriate file names (without quotes). Since this file should be in the same directory as programs, you will not need any special paths, though CMake will tell you if there are no files with that name.

Note

If you did not label any of your HIP files with .hip, make a copy of them with the .hip extension instead of a C++ extension. A later sub-step will handle dealing with HIP files that do not have a .hip extension.

cmake_minimum_required(VERSION 3.21)

# Enable both CXX and HIP languages
project(BootcampBuilds LANGUAGES HIP CXX)

# Locate the MPI library
find_package(MPI REQUIRED)

# Locate the RCCL library
find_package(RCCL REQUIRED)

if(NOT CMAKE_HIP_ARCHITECTURES)
    message(FATAL "No CMAKE_HIP_ARCHITECTURES was set")
endif()

# Step 0: C++ Program
# ==============================================================================
add_executable(step0 <ADD ME!>)

# Step 1: MPI Program
# ==============================================================================
add_executable(step1 <ADD ME!>)
target_link_libraries(step1 PRIVATE MPI::MPI_CXX)

# Step 2: HIP Program
# ==============================================================================
add_executable(step2 <ADD ME!>)

# Step 3: MPI + HIP Program
# ==============================================================================
add_executable(step3 <ADD ME!>)
target_link_libraries(step3 PRIVATE MPI::MPI_CXX)
# target_compile_definitions(step3 PUBLIC TEST_GPU_AWARE)

# Step 4: RCCL Program
# ==============================================================================
add_executable(step4 <ADD ME!>)
target_link_libraries(step4 PRIVATE roc::rccl)

# Step 5: RCCL + MPI Program
# ==============================================================================
add_executable(step5 <ADD ME!>)
target_link_libraries(step5 PRIVATE roc::rccl MPI::MPI_CXX)

Once you have the CMakeLists.txt filled out, do mkdir build and then cd build. This is the build directory where CMake will build your executables. Before proceeding to the next part of the build process, go ahead and get on a fresh instance of your terminal session. Since specific modules can influence things in the environment, and therefore what CMake sees, it is always a good idea to make sure your environment is fresh when trying to debug issues. As such, the only module that needs to be loaded for the first part of this step is a "rocm" module (i.e., module load rocm).

After loading that module, do the command cmake -DCMAKE_HIP_ARCHITECTURES=gfx942 .. inside the build folder. If all went according to plan, you should see CMake successfully configure itself to be ready to build the project. Once done, you can do make to see it build the executables. Finally, try running to double check that they all work!

Note

You may see other example of CMake compilations with something similar to CC=mpic++ cmake .. or cmake -DCMAKE_CXX_COMPILER=mpic++. Both of these examples are explicitly telling CMake that it should use mpic++ as the C++ compiler, and in most cases, this is often enough to help get CMake over some linking issues. Similarly, some Cray documentation recommends that you do -DCMAKE_C_COMPILER=cc -DCMAKE_CXX_COMPILER=CC when you compile. However, like most things, this is fine until it breaks (such as when HIP is introduced into the build). The rest of the step focuses more on what is happening under the hood, and assumes none of these techniques have been used unless otherwise stated.

Step 7b: CMake and Cray MPICH GTL Library

Unfortunately, while CMake is very good at figuring out all of the dependencies different MPI implementations have, there are still occasional corners where something isn't quite right. The perfect example of this is finding the GTL library GPU-aware MPI programs need when running with cray-mpich. To verify that CMake does not inherently find this dependency, remove the comment (the #) in front of the target_compile_definitions in the Step 3 build. This line of CMake code will add the -DTEST_GPU_AWARE flag you used back in Step 3. Remove all of the file in the build folder, re-run the CMake command from 7a, run make, and then take a look at what libraries are linked in the Step 3 executable (use ldd). Assuming there is nothing in your environment (and as of August 2026), you should see that there is no GTL library in this program.

To resolve this issue, there are a handful of options:
1. Try out different C/C++ compilers (such as Cray's cc or CC wrappers).
2. On LLNL systems, you can load a module that has a magic wrapper that tweaks how MPI programs are compiled. For example, doing module load rocmcc instead of (or in addition to) module load rocm will cause the GTL library to be linked for ALL libraries or executables linked to MPI
3. A more explicit, but still "secret knowledge" approach involves adjusting certain CMake flags when configuring the project. Specifically, Cray provides the variables PE_MPICH_GTL_DIR_<GPU_ARCH> and PE_MPICH_GTL_LIBS_<GPU_ARCH> that hold the specific values one should pass to -L and -l, respectively. Note that "" is the GPU architecture of the system you are on. For Tuolumne, this would be "amd_gfx942". Combined with CMake's ability to pass compiler and linker flags at command line, you can combine these two to tell CMake that there are more flags it needs to bring in:

# For projects that are building a library 
cmake <other flags> "-DCMAKE_SHARED_LINKER_FLAGS=${PE_MPICH_GTL_DIR_amd_gfx942} ${PE_MPICH_GTL_LIBS_amd_gfx942}" 
# For projects that are building executables
cmake <other flags> "-DCMAKE_EXE_LINKER_FLAGS=${PE_MPICH_GTL_DIR_amd_gfx942} ${PE_MPICH_GTL_LIBS_amd_gfx942}" ..
# If project does both, add both CMake flags from above.

4. Tweak your CMake code to explicitly search for the GTL library. This option is left up to the user to try out, as it generally the least portable since you have to add Cray-specific logic into your CMake code. The CMake code uses environment checks on the "PE_MPICH_GTL" variables above to combine them with a targeted find_library CMake call to get the necessary library included.

None of these options are the "best". In practice Option 2 is generally the "easiest" to use with Option 3 as a fallback when things get complicated. However, it is always a good idea to check what CMake found by using ldd on the final outputs, or by using make VERBOSE=1 when building to see what libraries paths are being used.s

Tip

CMake cannot set environment variables for when you run your applications. You (and CMake) can set environment variables that will influence compilation, but setting any flags that influence runtime behaviors (i.e. MPICH_GPU_SUPPORT_ENABLED), will not change anything in your build.

Step 7c: C++ HIP Files

The last part of this step is exploring what happens when you don't use a .hip file extension for code that is using only HIP libraries and for code that contains GPU code. For this example, make a copy of the file used for Step 5 and Step 3, and name them step7cA.cpp and step7cB.cpp, respectively (watching out for that Step 5 needs to be the "A" file here).

To do this, add the following CMake code to the bottom of you CMakeLists.txt:

# Step 7c: C++ HIP program
# ==============================================================================
add_executable(step7cA step7cA.cpp)
target_link_libraries(step7cA PRIVATE roc::rccl MPI::MPI_CXX)
add_executable(step7cB step7cB.cpp)
target_compile_definitions(step7cB PUBLIC TEST_GPU_AWARE)
target_link_libraries(step7cB PRIVATE MPI::MPI_CXX)

After re-running the CMake command and start rebuilding the files, you should see that only the first added file successfully builds. This is because this example has no actual GPU code in it, only HIP, RCCL, and MPI library calls. Taking a closer look at the error

/step7cB.cpp:3:10: fatal error: 'hip/hip_runtime.h' file not found
    3 | #include <hip/hip_runtime.h>

along with the specific targets used for its target_link_libraries command, the first issue is that there CMake was not told that this file needs to link to the HIP library. And the keen will notice that technically the first file also probably needs this linkage (though roc::rccl likely implies this dependency somewhere down the line). To resolve that, add the following line (up near the other find_package commands)

find_package(hip REQUIRED)

Then, down with the target_link_libraries for each target, add hip::host after the PRIVATE keyword, like so:

target_link_libraries(step7cA PRIVATE hip::host roc::rccl MPI::MPI_CXX)
target_link_libraries(step7cB PRIVATE hip::host MPI::MPI_CXX)

Now if both examples only needed to link to MPI and HIP libraries, the code would be done. However, because the second file contains actual GPU code, using a standard C++ compiler will get upset since it is not C++ syntax:

step7cB.cpp:64:21: error: expected expression
   64 |     initDataKernel<<<blocks, threadsPerBlock>>>(d_sendbuf, rank, buffer_size);
      |                     ^
step7cB.cpp:64:47: error: expected expression
   64 |     initDataKernel<<<blocks, threadsPerBlock>>>(d_sendbuf, rank, buffer_size);
      |                                               ^

To get around this, one can tell CMake that file has a particular language with the following tweak (anywhere in the file is fine):

set_property( SOURCE step7cB.cpp PROPERTY LANGUAGE HIP)

Re-run CMake (after cleaning out the build directory), rebuild the files, and you should see that it all builds successfully!

Tip

While changing the language of a given file in CMake is valid code, it is generally a better idea to have the files with GPU code have the appropriate file extension (.hip or .cu).

Step 7d: CMake Language Ordering

One final consideration about using CMake relates to which compilers it uses. In particular, the order of the languages in the project() line can affect which compiler is being used:

# Output from "project(BootcampBuilds LANGUAGES HIP CXX)"
-- The HIP compiler identification is Clang 19.0.0
-- The CXX compiler identification is Clang 19.0.0
-- Check for working HIP compiler: /opt/rocm-6.4.3/lib/llvm/bin/clang++ - skipped
-- Check for working CXX compiler: /opt/rocm-6.4.3/lib/llvm/bin/clang++ - skipped
-- Found MPI_CXX: /opt/cray/pe/mpich/9.0.1/ofi/cray/20.0/lib/libmpi_cray.so (found version "4.0")
-- Found MPI: TRUE (found version "4.0")
# Output from "project(BootcampBuilds LANGUAGES CXX HIP)"
-- The CXX compiler identification is GNU 8.5.0
-- The HIP compiler identification is Clang 19.0.0
-- Check for working CXX compiler: /usr/lib64/ccache/c++ - skipped
-- Check for working HIP compiler: /opt/rocm-6.4.3/lib/llvm/bin/clang++ - skipped
-- Found MPI_CXX: /opt/cray/pe/mpich/9.0.1/ofi/cray/20.0/lib/libmpi_cray.so (found version "4.0")
-- Found MPI: TRUE (found version "4.0")

Notice how a different CXX compiler was chosen in the second case. While this ultimately does not affect the compilation of any of these tests in this example, paying attention to the outputs here can give key insights into what might be off when things error out. Feel free to play around with giving different compilers to CMake (or by loading different modules) and see how/if they influence CMake!

Note

While the GNU compilers were initially discouraged from use in Step 0, CMake is smart enough to link to the correct libraries underneath so that the appropriate code is called, which is likely why the programs work without issues. However, this particular example likely should try and have Clang-based compilers for any language to avoid any weird potential edge cases.

Conclusion

This module covered the basics of compiling MPI and GPU programs, albeit slightly tailored to compilation on LLNL's Tuolumne. The general strategy of checking modules, compilers, libraries, and environment flags is a good approach to help sort out compilation and basic GPU errors when running MPI applications. If there are any other suggestions, or any of the advice seems out of date, feel free to let us know!