OpenCL Matrix-Vektor-Multiplikation

[Rüdiger Knörig (Gast)]

Zum Erlernen von OpenCL habe ich mir eine Vektor-Matrix-Multiplikation 
geschrieben, bei der ich mir aber nicht sicher bin, inwieweit die 
dahinterstehenden Überlegungen richtig sind.

Prinzipiell verstehe ich Workgroups so, daß mit diesen das Gesamtproblem 
in von dem jeweiligen Gerät gleichzeitig ausführbare Häppchen zerlegt 
wird.

Deshalb habe ich die elementweise Zeilenvektor/Vektormultiplikation 
durch je eine Workgroup ausführen lassen, wobei erst einmal die erste 
Einheit der Workgroup den Produktvektor aufsummiert.

Somit wäre die "Global Size" die Größe der Matrix und die "Local Size" 
die Anzahl der Zeilenvektoren / Produktvektorgröße.

Kann ich generell (bei mir funktioniert das) den Kernel dann so 
parameterisieren:

 cl::NDRange globalSize(A.getNumberOfRows()*A.getNumberOfColumns());
 cl::NDRange localSize(A.getNumberOfColumns());
 rcpp::KernelFunctor vektorSquareFunc(matrixMultiplicationKernel,queue,globalSize,localSize);
        vektorSquareFunc(cl::Buffer(A),
                         cl::Buffer(x),
                         cl::Buffer(y),
                         A.getNumberOfRows(),
                         A.getNumberOfColumns(),
                         rcpp::KernelFunctor::createLocalMemoryBuffer<double>(A.getNumberOfColumns()),
                         rcpp::KernelFunctor::createLocalMemoryBuffer<double>(A.getNumberOfColumns()));

Den Kernel selbst habe ich so geschrieben, wobei mich momentan der 
Verdacht überkommt, daß das "Cachen" des Produktvektors in einen lokalen 
Puffer unnötig ist - es sei denn, die physische Arbeitsgruppe wird 
wieder rangenommen:

/**
 * @brief Matrix multiplication kernel.
 * @param A The matrix as concencated row vectors.
 * @param x Multiplicant vector (size: number of columns in A)
 * @param y Result vector A*x (size: number of rows in A)
 * @param nRows Number of rows in the matrix.
 * @param nCols Number of columns in the matrix.
 * @param xbuffer Local memory buffer for the multiplicant vector.
 */

 

__kernel void matrixMultiplication(__global double *A,__global double *x,__global double *y,unsigned int nRows,unsigned int nCols,__local double *xbuffer,__local double *resultBuffer)
{
    int column=get_local_id(0); // equals the column number
    int row=get_group_id(0); // fetch the workgroup number - should be the row number
    int globalID=get_global_id(0);

    xbuffer[column]=x[column];

    barrier(CLK_LOCAL_MEM_FENCE);

    resultBuffer[column]=xbuffer[column]*A[globalID];

    barrier(CLK_LOCAL_MEM_FENCE);

    double sum=0;

    if(column==0)
    {
        for(int k=0;k<nCols;k++) sum += resultBuffer[k];
        y[row]=sum;
    }

}

[Olga (Gast)]

Hm, welche obskure Verschwörungstheorie möchtest du denn damit beweisen?