La idea de mi programa simple que he estado tratando de escribir es tomar información del usuario para ver qué tan grande es la matriz para multiplicar.

Estoy buscando tomar la entrada x por x, actualmente no estoy buscando multiplicar dos tamaños diferentes en este momento.

¿Cómo sugerirían ustedes que haga esto?

Lo siento, mi pregunta no fue lo suficientemente clara, quiero modificar este kernel para que pueda manejar una matriz de cualquier tamaño (donde xey son equivalentes para mantenerlo simple). En lugar de múltiplos de 16.

No estoy seguro de si necesitaría mi código actual, pero aquí está el código del kernel:

// CUDA Kernel
__global__ void matrixMul( float* C, float* A, float* B, int wA, int wB,size_t block_size)
{
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    int aBegin = wA * block_size * by;
    int aEnd   = aBegin + wA - 1;
    int aStep  = block_size;

    int bBegin = block_size * bx;

    int bStep  = block_size * wB;
    float Csub=0;

    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
    {
        extern __shared__ float As[];
        extern __shared__ float Bs[];
        extern __shared__ float smem[];

        smem[ty*block_size+tx] = A[a + wA * ty + tx];

        smem[block_size*block_size+ty*block_size+tx]  = B[b + wB * ty + tx];

        __syncthreads();

        for (int k = 0; k < block_size; ++k)
            Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ;

        __syncthreads();
    }

    int c = wB * block_size * by + block_size * bx;
    C[c + wB * ty + tx] = Csub;


}

Actualización: decidí ir con el relleno de cero. Sin embargo, obtengo respuestas incorrectas. Tome la matriz A 2x2, acolchada a 16x16:

5.000 0.000 9.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000

Matrix B, 2x2 acolchado a 16x16:

7.000 4.000 8.000 7.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000

Entonces, el resultado de C que obtengo es correcto:

35.000 20.000 40.000 35.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000

Sin embargo, si quita los ceros, la matriz debería ser: A:

5.000 0.000
9.000 0.000

SI:

7.000 4.000
8.000 7.000

C debe ser:

35.000 20.000
63.000 36.000

Sin embargo, las dos matrices C no son iguales.

2
Dan 12 feb. 2012 a las 21:25
¿Cuál es tu pregunta? ¿Está preguntando cómo obtener información del usuario?
 – 
harrism
13 feb. 2012 a las 08:09
Si entendí correctamente sus preguntas anteriores, su pregunta real es cómo modificar este código del kernel (en sí mismo una versión muy ligeramente modificada del ejemplo de multiplicación de matrices CUDA SDK) para que pueda usarse para multiplicar matrices arbitrarias. tamaño, a diferencia de los múltiplos redondos del tamaño del bloque del núcleo. ¿Podrías editar tu pregunta para reflejar esto? Por el momento, no está muy claro qué es lo que realmente está preguntando.
 – 
talonmies
13 feb. 2012 a las 14:37
@talonmies, tienes razón. Eso es exactamente lo que estoy buscando
 – 
Dan
13 feb. 2012 a las 18:47
Hola Dan, ¿encontraste el código correcto para tu implementación arriba?
 – 
user990246
18 abr. 2012 a las 00:45
Lo hice gracias. Pero creo que está limitado a 16 bloques. Solo lo estoy revisando para ver si puedo cambiar eso.
 – 
Dan
19 abr. 2012 a las 07:45

1 respuesta

La mejor respuesta

Esta no es una pregunta muy clara, por lo que esta respuesta es una especie de conjetura basada en lo que ha preguntado anteriormente en varias preguntas bastante similares.

Un buen punto de partida para entender cómo hacer este tipo de operación es volver al principio y pensar en el problema de la multiplicación matriz-matriz desde los primeros principios. Está interesado en el código para calcular el producto escalar de dos matrices, C = AB . La restricción que tiene es que el núcleo que está utilizando solo puede calcular productos de matrices que son múltiplos redondos de algún tamaño de bloque interno. ¿Entonces que puedes hacer?

Una forma de ver el problema es imaginar que las matrices A y B eran matrices de bloques . La multiplicación de matrices se puede escribir así:

enter image description here

Y la matriz resultante C puede formarse mediante combinaciones de los productos de las ocho submatrices en A y B :

enter image description here

Puede que no sea inmediatamente obvio cómo esto ayuda a resolver el problema, pero consideremos un ejemplo concreto:

  1. Tiene un núcleo de multiplicación de matrices óptimo que usa un tamaño de bloque interno de 32, y solo es correcto cuando las matrices son múltiplos redondos de ese tamaño de bloque.
  2. Tienes un par de matrices cuadradas de 1000x1000 para multiplicar.

Estos primeros hechos implican que su kernel solo puede resolver correctamente un producto de 1024x1024 o un producto de 992x992, pero no la operación de 1000x1000 que necesita.

Si decide utilizar un producto de 1024x1024, puede utilizar la idea de descomposición de bloques para formular el problema de la siguiente manera:

enter image description here

Donde O nn denota una matriz de ceros de tamaño adecuado. Ahora tiene un par de matrices 1024x1024 y su producto dará como resultado

enter image description here

Es decir. el bloque superior izquierdo es una matriz de 1000x1000 que contiene AB . Esto es efectivamente relleno de cero para lograr el resultado correcto. En este ejemplo, significa que se realiza aproximadamente un 7% más de cálculo de lo necesario. Si eso es importante o no, probablemente sea específico de la aplicación.

El segundo enfoque sería usar el kernel básico para calcular un producto de 992x992, luego elaborar una estrategia para lidiar con los otros siete productos en la versión descompuesta en bloque del cálculo, algo como esto:

enter image description here

Siendo A 11 y B 11 matrices de 992x992, y O nn son matrices cero como antes. A primera vista, esto no parece muy útil, pero vale la pena recordar que todos los cálculos para hacer la matriz del lado derecho contienen solo alrededor del 1.2% de los cálculos totales requeridos para calcular el producto de la matriz. Podrían realizarse fácilmente en la CPU del host mientras la GPU realiza el cálculo principal y luego agregarse al resultado de la GPU para formar la matriz final. Debido a que la API de CUDA es asincrónica, la mayor parte de ese cálculo de host se puede ocultar por completo y es efectivamente gratuito.

Esta respuesta contiene dos estrategias para hacer lo que está pidiendo sin cambiar más de una línea de su código de kernel actual . Obviamente, hay una tercera forma, que es modificar más radicalmente el núcleo en sí, pero eso es algo que debe probar usted mismo primero y luego pedir ayuda si su solución no funciona.

7
talonmies 13 feb. 2012 a las 18:18
Gracias por la explicación bien formada. Voy con la solución acolchada por el bien del tiempo. Tengo curiosidad acerca de dónde se le ocurrió un "7% más de cálculo".
 – 
Dan
14 feb. 2012 a las 02:37
Intenté ponerle relleno a cero, sin embargo, arroja resultados incorrectos. Pasé un tiempo tratando de resolverlo, pero parece que se encuentra en los 0 durante el relleno. Consulte mi OP editado para la actualización.
 – 
Dan
14 feb. 2012 a las 08:30
La forma en que ha realizado el relleno de ceros en su ejemplo es incorrecta. Y la cifra del 7% proviene de la diferencia entre 2 * 1024 ^ 3 y 2 * 1000 ^ 3, que son los recuentos de operaciones de los productos escalares en los dos tamaños.
 – 
talonmies
14 feb. 2012 a las 08:59
Hola talonmies, gracias por la gran respuesta. Mencionaste que la estrategia necesitaría una línea de cambio de código. Entonces, no se trata de crear una matriz dentro de una matriz, sino de tener una declaración if para verificar los límites, ¿correcto? Dado el código de Dan anterior, ¿cómo averiguaría los límites? Gracias
 – 
user990246
17 abr. 2012 a las 22:01