Как правильно распределить общую память в PyOpenCL с ядром, используя 2 общие переменные? - PullRequest
Новая
       57

Как правильно распределить общую память в PyOpenCL с ядром, используя 2 общие переменные?

0 голосов
/ 28 октября 2019

Я работаю с этот код реализация умножения матриц с использованием PyOpenCL. Я считаю, что при использовании этого ядра (при условии, что я) неправильно распределяю общую память рабочей группы.

Две переменные в этом ядре создаются из локальной памяти, должен ли параметр локальной памяти в вызове функции ядра в Python быть размером одной переменной или обеих?

У меня естьскорректировал размеры, размер мозаики TS и размеры входного массива и не смог получить правильный вывод. 

# class and kernel string - there should be a block indent here, sorry about this
Class MatrixMultiply:
#define TS 1
kernel = '''

// Tiled and coalesced version
__kernel void myGEMM2(const int M, const int N, const int K,
                  const __global float* A,
                  const __global float* B,
                  __global float* C) {

// Thread identifiers
const int row = get_local_id(0); // Local row ID (max: TS)
const int col = get_local_id(1); // Local col ID (max: TS)
const int globalRow = TS*get_group_id(0) + row; // Row ID of C (0..M)
const int globalCol = TS*get_group_id(1) + col; // Col ID of C (0..N)

// Local memory to fit a tile of TS*TS elements of A and B
__local float Asub[TS][TS];
__local float Bsub[TS][TS];

// Initialise the accumulation register
float acc = 0.0f;

// Loop over all tiles
const int numTiles = K/TS;
for (int t=0; t<numTiles; t++) {

    // Load one tile of A and B into local memory
    const int tiledRow = TS*t + row;
    const int tiledCol = TS*t + col;
    Asub[col][row] = A[tiledCol*M + globalRow];
    Bsub[col][row] = B[globalCol*K + tiledRow];

    // Synchronise to make sure the tile is loaded
    barrier(CLK_LOCAL_MEM_FENCE);

    // Perform the computation for a single tile
    for (int k=0; k<TS; k++) {
        acc += Asub[k][row] * Bsub[col][k];
    }

    // Synchronise before loading the next tile
    barrier(CLK_LOCAL_MEM_FENCE);
}

// Store the final result in C
C[globalCol*M + globalRow] = acc;
}
'''

def matrix_mul_optimized(self, a, b):
    dim_0 = a.shape[0]
    dim_1 = b.shape[1]
    common_dim = a.shape[1]   

    a = np.array(a, dtype=np.float32)  
    b = np.array(b, dtype=np.float32)  
    dummy_array = np.ones(shape=(dim_0,dim_1), dtype=np.float32) 

    mf = cl.mem_flags
    a_buf = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a)
    b_buf = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b)
    c_buf = cl.Buffer(self.ctx, mf.WRITE_ONLY, dummy_array.nbytes)


    prg = cl.Program(self.ctx, self.kernel_optimized).build()

    # HERE IS WHERE I AM CONCERNED!! Tile size 'TS' is 1, should the local memory be 1? Or 2 because the kernel has two variables of this size?
    prg.MatrixMul(self.queue, dummy_array.shape, dummy_array.shape, 
    np.int32(dim_0), 
    np.int32(dim_1),
    np.int32(common_dim),
    a_buf, b_buf, c_buf)


    product = np.zeros((dim_0,dim_1), dtype=np.float32)
    cl.enqueue_copy(self.queue, product, c_buf)

    return product.astype(int)

main 

if __name__ == '__main__':

np.random.seed(7)
# Array dims
M = 2
N = 2

multiplier = MatrixMultiply()

a = np.random.randint(1,5,(M,N)).astype(np.float32)
b = np.random.randint(1,5,(N,M)).astype(np.float32)

matrix_mul_optimized = multiplier.matrix_mul_optimized(a, b)

Я ожидаю правильного умножения матрицОднако я получаю случайное число и разные числа в массиве возврата 2x2 каждый раз, когда я изменяю такие параметры, как размер локальной памяти. Ядро взято прямо из учебника, поэтому я могу только предположить, что что-то в моей оболочке Python неверно.

...