SLI for multiple GPUs

Question

I am new to CUDA programming, and I am working on a problem that requires multiple GPUs in one machine. I understand that for better graphics programming multiple GPUs need to be combined via SLI. However, for CUDA programming do I need to combine GPUs via SLI as well?

Answer 1

You can use CUDA on multiple GPUs without SLI and even between different architectures of CUDA but you have to write extra codes to divide the work and synchronize the divided sub-works. Here is a simple program doing load-balancing on 3 GPUs for the sample kernel vectorAdd (GT1030 GPU a Pascal architecture GPU + two K420 GPUs that are Kepler architecture, working without a problem together for same task pool):

/**
 * Copyright 1993-2015 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 *
 */

/**
 * Vector addition: C = A + B.
 *
 * This sample is a very basic sample that implements element by element
 * vector addition. It is the same as the sample illustrating Chapter 2
 * of the programming guide with some additions like error checking.
 */

#include <stdio.h>

// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>

#include <helper_cuda.h>

// for load balancing between 3 different GPUs
#include "LoadBalancerX.h"

/**
 * CUDA Kernel Device code
 *
 * Computes the vector addition of A and B into C. The 3 vectors have the same
 * number of elements numElements.
 */
__global__ void
vectorAdd(const float *A, const float *B, float *C, int numElements)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numElements)
    {
        C[i] = A[i] + B[i];
    }
}


#include<iostream>
#include<map>
int
main(void)
{

    int numElements = 1500000;
    int numElementsPerGrain = 50000;
    size_t size = numElements * sizeof(float);

    float *h_A; cudaMallocHost((void**)&h_A,size);
    float *h_B; cudaMallocHost((void**)&h_B,size);
    float *h_C; cudaMallocHost((void**)&h_C,size);


    for (int i = 0; i < numElements; ++i)
    {
        h_A[i] = rand()/(float)RAND_MAX;
        h_B[i] = rand()/(float)RAND_MAX;
    }



    /*
     * default tutorial vecAdd logic

    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);


    int threadsPerBlock = 256;
    int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
    cudaGetLastError();


    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    */

    /* load-balanced 3-GPU version setup */
    class GrainState
    {
    public:
        int offset;
        int range;
        std::map<int,float *> d_A;
        std::map<int,float *> d_B;
        std::map<int,float *> d_C;
        std::map<int,cudaStream_t> stream;
        ~GrainState(){
            for(auto a:d_A)
                cudaFree(a.second);
            for(auto b:d_B)
                cudaFree(b.second);
            for(auto c:d_C)
                cudaFree(c.second);
            for(auto s:stream)
                cudaStreamDestroy(s.second);
        }
    };

    class DeviceState
    {
    public:
        int gpuId;
        int amIgpu;
    };

    LoadBalanceLib::LoadBalancerX<DeviceState,GrainState> lb;
    lb.addDevice(LoadBalanceLib::ComputeDevice<DeviceState>({0,1})); // 1st cuda gpu in computer
    lb.addDevice(LoadBalanceLib::ComputeDevice<DeviceState>({1,1})); // 2nd cuda gpu in computer
    lb.addDevice(LoadBalanceLib::ComputeDevice<DeviceState>({2,1})); // 3rd cuda gpu in computer
    //lb.addDevice(LoadBalanceLib::ComputeDevice<DeviceState>({3,0})); // CPU single core

    for(int i=0;i<numElements;i+=numElementsPerGrain)
    {
        lb.addWork(LoadBalanceLib::GrainOfWork<DeviceState,GrainState>(
                [&,i](DeviceState gpu, GrainState& grain){
                    if(gpu.amIgpu)
                    {
                        cudaSetDevice(gpu.gpuId);
                        cudaStreamCreate(&grain.stream[gpu.gpuId]);
                        cudaMalloc((void **)&grain.d_A[gpu.gpuId], numElementsPerGrain*sizeof(float));
                        cudaMalloc((void **)&grain.d_B[gpu.gpuId], numElementsPerGrain*sizeof(float));
                        cudaMalloc((void **)&grain.d_C[gpu.gpuId], numElementsPerGrain*sizeof(float));
                    }
                },
                [&,i](DeviceState gpu, GrainState& grain){
                    if(gpu.amIgpu)
                    {
                        cudaSetDevice(gpu.gpuId);
                        cudaMemcpyAsync(grain.d_A[gpu.gpuId], h_A+i, numElementsPerGrain*sizeof(float), cudaMemcpyHostToDevice,grain.stream[gpu.gpuId]);
                        cudaMemcpyAsync(grain.d_B[gpu.gpuId], h_B+i, numElementsPerGrain*sizeof(float), cudaMemcpyHostToDevice,grain.stream[gpu.gpuId]);
                    }
                },
                [&,i](DeviceState gpu, GrainState& grain){
                    if(gpu.amIgpu)
                    {
                        int threadsPerBlock = 1000;
                        int blocksPerGrid =numElementsPerGrain/1000;
                        vectorAdd<<<blocksPerGrid, threadsPerBlock, 0, grain.stream[gpu.gpuId]>>>(grain.d_A[gpu.gpuId], grain.d_B[gpu.gpuId], grain.d_C[gpu.gpuId], numElements-i);
                    }
                    else
                    {
                        for(int j=0;j<numElementsPerGrain;j++)
                        {
                            const int index = j+i;
                            h_C[index]=h_A[index]+h_B[index];
                        }
                    }
                },
                [&,i](DeviceState gpu, GrainState& grain){
                    if(gpu.amIgpu)
                    {
                       cudaMemcpyAsync(h_C+i, grain.d_C[gpu.gpuId], numElementsPerGrain*sizeof(float), cudaMemcpyDeviceToHost,grain.stream[gpu.gpuId]);
                    }
                },
                [&,i](DeviceState gpu, GrainState& grain){
                    if(gpu.amIgpu)
                    {
                        cudaStreamSynchronize(grain.stream[gpu.gpuId]);
                    }
                }
        ));
    }

    /* load-balance setup end*/

    // run 100 times
    size_t nanoseconds=0;

    for(int i=0;i<100;i++)
    {
        nanoseconds += lb.run();

    }

    std::cout<<nanoseconds/100.0<<"ns  ("<<((numElements*12.0/(nanoseconds/100.0)))<<"GB/s)"<<std::endl;


    std::cout<<"??"<<std::endl;

    for (int i = 0; i < numElements; i+=numElementsPerGrain)
    {
        std::cout<<h_A[i]<<" + "<<h_B[i]<<" = "<<h_C[i]<<std::endl;
    }
    auto z = lb.getRelativePerformancesOfDevices();
    std::cout<<"work distribution to devices:"<<std::endl;
    for(auto zz:z)
    {
        std::cout<<zz<<"% ";
    }
    std::cout<<std::endl;
    cudaFreeHost(h_A);
    cudaFreeHost(h_B);
    cudaFreeHost(h_C);

    return 0;
}

Answer 2

No, in general you don't want to use SLI if you plan on using the GPUs for compute instead of pure graphics applications. You will be able to access both GPUs as discrete devices from within your CUDA program. Note that you will need to explicitly divide work between the GPUs.

I don't have an explanation for why SLI isn't desirable for compute applications, but it's what I've read on the Nvidia forums and heard from others in IRC channels.