5
Unten ist ein opencl-Kernel, der blockierte Matrixmultiplikation für mehrere unabhängige Matrizen durchführt. selectMatrixA und selectMatrixB speichern mehrere Matrizen (gleiche Größe und quadratische Matrizen) in der Reihenfolge der Zeilen.Optimierung der Stapelmatrixmultiplikation opencl code
// Matrix multiplication: C = A * B.
#define BLOCK_SIZE 20
#define MATRIX_SIZE 100 * 100
#define BLOCK_DIMX 5 // Number of blocks in the x dimension
__kernel void
batchedMatrixMul(__global float *selectMatrixC, __global float *selectMatrixA, __global
float *selectMatrixB, int wA, int wB)
{
// Block index
int bx = get_group_id(0);
int by = get_group_id(1);
__global float *C = selectMatrixC + (bx/BLOCK_DIMX) * MATRIX_SIZE;
__global float *A = selectMatrixA + (bx/BLOCK_DIMX) * MATRIX_SIZE;
__global float *B = selectMatrixB + (bx/BLOCK_DIMX) * MATRIX_SIZE;
int tx = get_local_id(0);
int ty = get_local_id(1);
float Csub = 0;
// Identify the row and column of the C matrix to work on
int Row = (by * BLOCK_SIZE) + ty;
int Col = ((bx %(BLOCK_DIMX)) * BLOCK_SIZE) + tx;
// Declaration of the local memory array As used to store the sub-matrix of A
__local float As[BLOCK_SIZE][BLOCK_SIZE];
// Declaration of the local memory array Bs used to store the sub-matrix of B
__local float Bs[BLOCK_SIZE][BLOCK_SIZE];
// Loop over all the sub-matrices of A and B required to compute the block sub-matrix
for (int m = 0; m < wA/BLOCK_SIZE; ++m)
{
// Load the matrices from global memory to local memory. Each thread loads one
//element of each matrix
As[ty][tx] = A[Row * wA + m * BLOCK_SIZE + tx];
Bs[ty][tx] = B[(m * BLOCK_SIZE + ty)*wA + Col];
// Synchronize to make sure the matrices are loaded
barrier(CLK_LOCAL_MEM_FENCE);
// Multiply the two matrices together each thread computes one element of the block
//sub-matrix
for (int k = 0; k < BLOCK_SIZE; ++k)
Csub += As[ty][k] * Bs[k][tx];
// Synchronize to make sure that the preceding computation is done before loading
//two new sub-matrices of A and B in the next iteration
barrier(CLK_LOCAL_MEM_FENCE);
}
// Write the block sub-matrix to device memory each thread writes one element
C[Row * wA + Col] = Csub;
}
Hier ist, wie ich den Kernel starten:
localWorkSize[0] = BLOCK_SIZE;
localWorkSize[1] = BLOCK_SIZE;
// for a 100 X 100 matrix, MATRIX_DIMX = MATRIX_DIMY = 100
globalWorkSize[0] = MATRIX_DIMX * NUM_MATRICES;
globalWorkSize[1] = MATRIX_DIMY ;
cl_event event;
errcode = clEnqueueNDRangeKernel(clCommandQueue,
clKernel, 2, NULL, globalWorkSize,
localWorkSize, 0, NULL, &event);
Im Folgenden sind einige Zahlen Leistung, wenn diese auf einem NVIDIA Grid K520 läuft:
1. matrix size:100 X 100 . Number of matrices = 20000. Time taken for multiplication =
0.262 seconds. As shown in the code, the block size was set to 20. Block size of 10 was
slower. This calculates to around 152 GFLOPS
2. matrix size: 10000 X 10000. Number of matrices = 1. Time taken for multiplication = 10.6
seconds. Here also the block size was 20. Using a block size of 50 is not possible due to
the size of the local memory.
Kann mir bitte jemand helfen zu verstehen, Warum läuft der Code langsam und warum 2. ist so viel langsamer als 1. Ich bin neu in OpenCL und ich möchte lernen, wie man Code basierend auf den zugrunde liegenden architektonischen Details optimiert.
Haben Sie einen einfachen Dot-Produktansatz versucht, um das Problem zu lösen und eine Basislinien-Leistungsmessung zu erhalten? Vielleicht sollte jedes Arbeitselement ein einzelnes Element in C berechnen, indem eine 100-Punkt-100-Operation ausgeführt wird. – mfa
müssen Ihre Matrixgrößen 100x100 und 10k x 10k sein? – mfa
ist es eine Voraussetzung, eine vollständige Multiplikation mit einer einzigen Arbeitsgruppe zu lösen, oder kann dies mit mehreren Gruppen geschehen? – mfa