¿Se puede optimizar este código OpenCL?

Question

estoy trabajando en una pieza de código OpenCL para una función de matriz especializada: para un vector Dx1v, dos DxD matrices A y B y una constante c, volver 1xD vector r donde r[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])

A continuación se muestra lo que tengo hasta ahora, pero funciona anormalmente lento. Una versión sin sumar que devuelve una matriz DxD es aproximadamente diez veces más rápida. Se llama desde PyOpenCL si eso hace alguna diferencia.

¿Ha hecho algo mal? ¿Podría ser optimizado?

#define D 1000 
... 

    __kernel void element_mult(
     __global float *result, 
     __global const float *vector, 
     __global const float *matrix, 
     __global const float *matrix2, 
     const float factor) 
     { 
     int y = get_global_id(1); 
     float sum = 0; 
     for(int k = 0; k < D; k++) 
     { 
      sum += vector[k] * matrix[(y*D) + k] 
      * matrix2[(y*D) + k ]; 
     } 
     result[y] = sum * factor; 
     } 

Saludos!

Answer 1

Optimización # 1: make vector __local.

Mi primer paso en esto obtuvo una mejora decente en el rendimiento. Noté que cada vector [k] se lee un total de D veces, así que lo copié a un __local. Esto solo es posible porque D es lo suficientemente pequeño como para permitir esto. El kernel como lo tienes arriba sufre de una terrible relación ALU: fetch de 0.08 tanto en el 5870 como en el 6970 gpus. Incluso los gpus más lentos todavía están esperando el acceso a la memoria.

#define D 1000 
    __kernel void element_mult(
    __global float *result, 
    __global const float *vector, 
    __global const float *matrix, 
    __global const float *matrix2, 
    const float factor) 
    { 
     int y = get_global_id(0); 
     float sum = 0; 

     __local float vectCopy[D]; 
     int ls = get_local_size(0); 
     int lid = get_local_id(0); 
     for(int i=0;i<D;i+=ls){ 
      vectCopy[i+lid] = vector[i+lid]; 
     } 
     mem_fence(CLK_LOCAL_MEM_FENCE); 

     for(int k = 0; k < D; k++) 
     { 
      sum += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ]; 
     } 
     result[y] = sum * factor; 
    } 

Con este cambio, APP perfilador está mostrando una nueva ALU: ir a buscar proporción de 0,20 para el 5870 y el 6970 GPU. Los tiempos promedio cambiaron de 1513 -> 1034 y 1261 -> 861 en las mismas tarjetas. El gpus de gama baja ahora está vinculado por ALU en lugar de fetch. (Mayor de 4: 1 ratio)

Opimization # 2: calcular cada resultado [y] el uso de un grupo de trabajo entero.

Usted tendría que hacer esto Identificación del D eran mucho más grande (100k +). La idea es obtener el mejor patrón de acceso a la memoria utilizando el grupo de trabajo para calcular un solo elemento del resultado a la vez. Definí que ls (tamaño local) es 64 aquí, porque funciona en mi hardware, al igual que en la mayoría de los proveedores. El tamaño del grupo de trabajo que usa desde el lado del host tendrá que ser 64 a menos que cambie esa definición. Necesita ser definido para crear el almacenamiento de la suma [ls] como __local, y no me gusta pasar variables __local de tamaño variable a mis núcleos.

resultados: 5870 ALU: fetch = 0,59: 1, avg = 708. 6970 ALU: fetch = 0.72, avg = 590. De acuerdo con el perfilador de APP, esto es aproximadamente dos veces más rápido que su listado original.

#define D 1000 
#define ls 64 
__kernel void element_mult(
__global float *result, 
__global const float *vector, 
__global const float *matrix, 
__global const float *matrix2, 
const float factor) 
{ 
    __local float vectCopy[D]; 
    int lid = get_local_id(0); 
    for(int i=0;i<D;i+=ls){ 
     vectCopy[i+lid] = vector[i+lid]; 
    } 
    mem_fence(CLK_LOCAL_MEM_FENCE); 

    int ng = get_num_groups(0); 
    int gid = get_group_id(0); 
    int y, k; 
    __local float sum[ls]; 
    for(y = gid; y < D; y+=ng){ 
     for(k = lid; k < D; k+=ls) 
     { 
      sum[lid] += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ]; 
     } 
     if(lid==0){ 
      result[y] = sum[0]; 
      for(k=1;k<ls;k++){ 
       result[y] += sum[k]; 
      } 
      result[y] *= factor; 
     } 
     mem_fence(CLK_LOCAL_MEM_FENCE); 
    } 
} 

EDIT: APP perfilador = AMD APP KernelAnalyzer