Boost GPU Memory Performance with No Code Changes Using NVIDIA CUDA MPS

#include 
#include 
#include 

// Helper macro to ascertain for CUDA errors
#define CUDA_CHECK_FAILURE(x) 
if (cudaSuccess != (cudaError_t)x)
{
    const char* errName = cudaGetErrorName(x);
    const char* errStr = cudaGetErrorString(x);
    printf("%s:%d - %s: %sn", __FILE__, __LINE__, errName, errStr);
    exit(EXIT_FAILURE);
}

// Device memory variable to make use of to stop the compiler from optimizing away the memory access
__device__ volatile int dummy;

// Trivial kernel to the touch the memory so we are able to measure latency
__global__ void accessMemoryHighLatency(int *startAddress, size_t memorySizeInBytes) {
    for (int i = 0 ; i < memorySizeInBytes / sizeof(int) ; ++i) {
        dummy = atomicAdd(&startAddress[i], 1);
    }
}

// Function to launch the kernel and measure the runtime using CUDA events
float measureKernelRuntime(int *memoryDevPtr, size_t memorySizeInBytes, int numBlocks, int numThreads) {
    cudaEvent_t start = NULL, stop = NULL;
    float time = 0;

    CUDA_CHECK_FAILURE(cudaEventCreate(&start));
    CUDA_CHECK_FAILURE(cudaEventCreate(&stop));

    CUDA_CHECK_FAILURE(cudaEventRecord(start, 0));

    accessMemoryHighLatency<<>>(memoryDevPtr, memorySizeInBytes);
    CUDA_CHECK_FAILURE(cudaPeekAtLastError());

    CUDA_CHECK_FAILURE(cudaEventRecord(stop, 0));
    CUDA_CHECK_FAILURE(cudaEventSynchronize(stop));

    CUDA_CHECK_FAILURE(cudaEventElapsedTime(&time, start, stop));

    CUDA_CHECK_FAILURE(cudaEventDestroy(start));
    CUDA_CHECK_FAILURE(cudaEventDestroy(stop));

    return time;
}

int most important(int argc, char *argv[]) {
    size_t memorySizeInBytes = 32 * 1024 * 1024; // 32 MB
    int numBlocks = 32;
    int numThreads = 1;
    int numDevices = 0;
    float totalTime = 0;

    CUDA_CHECK_FAILURE(cudaGetDeviceCount(&numDevices));

    // Measure the runtime for every device
    for (int i = 0; i < numDevices; i++) {
        // Set the present device
        CUDA_CHECK_FAILURE(cudaSetDevice(i));
        
        // Allocate memory on the device
        int *memoryDevPtr;
        CUDA_CHECK_FAILURE(cudaMalloc(&memoryDevPtr, memorySizeInBytes));
        
        // Measure the runtime
        float time = measureKernelRuntime(memoryDevPtr, memorySizeInBytes, numBlocks, numThreads);
        totalTime += time;
        printf("Device %d - Total time: %f millisecondsn", i, time);
        
        // Free the memory
        CUDA_CHECK_FAILURE(cudaFree(memoryDevPtr));
    }

    printf("Average time: %f millisecondsn", totalTime / numDevices);

    return EXIT_SUCCESS;
}

$ nvidia-cuda-mps-control -d
# Optional step of explicitly creating an MPS server. This can also be done implicitly once we launch a CUDA application while MPS is energetic.
$ echo start_server -uid $UID | nvidia-cuda-mps-control
$ ./atomic_memory_access
Device 0 - Total time: 2320.550537 milliseconds
Device 1 - Total time: 2323.710938 milliseconds
Device 2 - Total time: 2334.533447 milliseconds
Device 3 - Total time: 2304.551025 milliseconds
Device 4 - Total time: 2304.328125 milliseconds
Device 5 - Total time: 2316.102295 milliseconds
Device 6 - Total time: 2306.165283 milliseconds
Device 7 - Total time: 2306.362061 milliseconds
Average time: 2314.537842 milliseconds

# Quit the MPS controller to cleanup the previous server.
$ echo quit | nvidia-cuda-mps-control
# Now repeat the above steps, with MLOPart enabled.
$ nvidia-cuda-mps-control -d
# Note that we must explicitly start the server with "-mlopart".
$ echo start_server -uid $UID -mlopart | nvidia-cuda-mps-control
$ ./atomic_memory_access
Device 0 - Total time: 1500.194946 milliseconds
Device 1 - Total time: 1475.914062 milliseconds
Device 2 - Total time: 1479.729492 milliseconds
Device 3 - Total time: 1480.196045 milliseconds
Device 4 - Total time: 1478.959106 milliseconds
Device 5 - Total time: 1490.808716 milliseconds
Device 6 - Total time: 1468.943237 milliseconds
Device 7 - Total time: 1479.297241 milliseconds
Device 8 - Total time: 1467.947632 milliseconds
Device 9 - Total time: 1476.900757 milliseconds
Device 10 - Total time: 1477.081421 milliseconds
Device 11 - Total time: 1490.295044 milliseconds
Device 12 - Total time: 1484.558594 milliseconds
Device 13 - Total time: 1481.660156 milliseconds
Device 14 - Total time: 1476.067383 milliseconds
Device 15 - Total time: 1484.143921 milliseconds
Average time: 1480.793457 milliseconds