2 video kartlı CUDA C programlama

Question

CUDA programlamada çok yeni ve nvidia tarafından sağlanan 'CUDA C Programlama Kılavuzu'nu okuyordum. (http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Programming_Guide.pdf)

Sayfa 25'te, matris çarpmasını yapan aşağıdaki C koduna sahiptir. Lütfen bu kodu iki cihazda nasıl çalıştırabilirim? (Bilgisayarımda yüklü iki nvida CUDA kapasiteli kart varsa). Bana bir örnek gösterebilir misin?

// Matrices are stored in row-major order: 
// M(row, col) = *(M.elements + row * M.stride + col) 
typedef struct { 
    int width; 
    int height; 
    int stride; 
    float* elements; 
} Matrix; 

// Get a matrix element 
__device__ float GetElement(const Matrix A, int row, int col) 
{ 
    return A.elements[row * A.stride + col]; 
} 

// Set a matrix element 
__device__ void SetElement(Matrix A, int row, int col, float value) 
{ 
    A.elements[row * A.stride + col] = value; 
} 

// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is 
// located col sub-matrices to the right and row sub-matrices down 
// from the upper-left corner of A 
__device__ Matrix GetSubMatrix(Matrix A, int row, int col) 
{ 
    Matrix Asub; 
    Asub.width = BLOCK_SIZE; 
    Asub.height = BLOCK_SIZE; 
    Asub.stride = A.stride; 
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col]; 
    return Asub; 
    } 

// Thread block size 
#define BLOCK_SIZE 16 

// Forward declaration of the matrix multiplication kernel 
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix); 

// Matrix multiplication - Host code 
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE 
void MatMul(const Matrix A, const Matrix B, Matrix C) 
{ 
    // Load A and B to device memory 
    Matrix d_A; 
    d_A.width = d_A.stride = A.width; d_A.height = A.height; 
    size_t size = A.width * A.height * sizeof(float); 
    cudaMalloc(&d_A.elements, size); 
    cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice); 
    Matrix d_B; 
    d_B.width = d_B.stride = B.width; d_B.height = B.height; 
    size = B.width * B.height * sizeof(float); 
    cudaMalloc(&d_B.elements, size); 
    cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice); 

    // Allocate C in device memory 
    Matrix d_C; 
    d_C.width = d_C.stride = C.width; d_C.height = C.height; 
    size = C.width * C.height * sizeof(float); 
    cudaMalloc(&d_C.elements, size); 

    // Invoke kernel 
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
    dim3 dimGrid(B.width/dimBlock.x, A.height/dimBlock.y); 
    MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 

    // Read C from device memory 
    cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost); 

    // Free device memory 
    cudaFree(d_A.elements); 
    cudaFree(d_B.elements); 
    cudaFree(d_C.elements); 
} 

// Matrix multiplication kernel called by MatMul() 
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) 
{ 
    // Block row and column 
    int blockRow = blockIdx.y; 
    int blockCol = blockIdx.x; 

    // Each thread block computes one sub-matrix Csub of C 
    Matrix Csub = GetSubMatrix(C, blockRow, blockCol); 

    // Each thread computes one element of Csub 
    // by accumulating results into Cvalue 
    float Cvalue = 0; 

    // Thread row and column within Csub 
    int row = threadIdx.y; 
    int col = threadIdx.x; 

    // Loop over all the sub-matrices of A and B that are 
    // required to compute Csub 
    // Multiply each pair of sub-matrices together 
    // and accumulate the results 
    for (int m = 0; m < (A.width/BLOCK_SIZE); ++m) 
    { 
     // Get sub-matrix Asub of A 
     Matrix Asub = GetSubMatrix(A, blockRow, m); 
     // Get sub-matrix Bsub of B 
     Matrix Bsub = GetSubMatrix(B, m, blockCol); 

     // Shared memory used to store Asub and Bsub respectively 
     __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 
     __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; 

     // Load Asub and Bsub from device memory to shared memory 
     // Each thread loads one element of each sub-matrix 
     As[row][col] = GetElement(Asub, row, col); 
     Bs[row][col] = GetElement(Bsub, row, col); 

     // Synchronize to make sure the sub-matrices are loaded 
     // before starting the computation 
     __syncthreads(); 

     // Multiply Asub and Bsub together 
     for (int e = 0; e < BLOCK_SIZE; ++e) 
      Cvalue += As[row][e] * Bs[e][col]; 

     // Synchronize to make sure that the preceding 
     // computation is done before loading two new 
     // sub-matrices of A and B in the next iteration 
     __syncthreads(); 
    } 

    // Write Csub to device memory 
    // Each thread writes one element 
    SetElement(Csub, row, col, Cvalue); 
} 

Answer 1

Birden fazla GPU'da bir CUDA çekirdeği çalıştırmak için "otomatik" bir yol yoktur.

Matris çarpma sorununu paralel olarak çalışabilen bağımsız işlemlere ayırmanın bir yolunu bulmanız gerekecektir (böylece her bir GPU'da paralel olarak). Basit bir örnek olarak:

C = A.BB1 ve B2 uygun B ve | columnwise concantenation belirtmektedir matrisin sütunları içeren boyutlu matrisleridir C = [A].[B1|B2] = [A.B1|A.B2] eşdeğerdir. Ayrı matris çarpma işlemleri olarak A.B1 ve A.B2 hesaplayabilir ve sonuçta elde edilen alt matrisleri ana belleğe geri kopyalarken birleştirme gerçekleştirebilirsiniz.

Uygun bir ayrıştırma düzenine sahip olduktan sonra, bunu CUDA 4.x API'sindeki standart çoklu-gpu olanaklarını kullanarak uygularsınız. CUDA API'lerini kullanarak çoklu GPU programlamaya genel bir bakış için, Paulius Micikevicius'ın bir video akışı ve PDF here olarak sunulan GTC 2012'den mükemmel konuşmasını izlemenizi tavsiye ederim.

Answer 2

temelleri Temel olarak, hangi mevcut konak iplik cudaSetDevice() arayarak çalışır GPU'yu ayarlayabilirsiniz CUDA C Programming Guide under section 3.2.6.

açıklanmıştır. Yine de, rutinlerinizi birden fazla GPU'ya bölmek için kendi kodunuzu yazmanız gerekir.