OpenCL如何改变cl_mem的内存地址？

Question

我想做一个子矩阵乘法。假设我有一个函数：

void MatMul(cl_mem A, cl_mem B, cl_mem C, int M, int K, int N)

其中A为M*K，B为K*N，C为M*N，A、B、C均为主机内存传过来的行主一维数组float *h_A, *h_B, *hC函数如下:

void ocl_push_array(cl_mem d_x, float *h_x, int n){
    size_t data_size = sizeof(float)*n;
    err = clEnqueueWriteBuffer(queue, d_x, CL_TRUE, 0, data_size, h_x, 0, NULL, NULL);
}

我想问:

如果我想进行子矩阵乘法，比如按行对 A 进行切片：

    // cl_mem A, B, C;
    for(int x=0; x<M; x+=16)
    {
        cl_mem A_sub = (cl_mem)((float *)A+x*K);
        cl_mem C_sub = (cl_mem)((float *)C+x*N);
        if((M-x+1)>=16)
            MatMul(A_sub, B, C_sub, 16, K, N);
        else
            MatMul(A_sub, B, C_sub, M-x+1, K, N);
    }

执行此操作的代码是否正确？我得到一个运行时间错误说："CL_INVALID_MEM_OBJECT" (-38) 当它将参数分配给 OpenCL 内核时 (clSetKernelArg).

我想做这个操作的原因是当我的输入矩阵A和B变大时，我发现矩阵乘法得到了错误的答案。

我的 OpenCL 内核是：

#define BLOCK_SIZE 16

#define AS(i, j) As[j + i * BLOCK_SIZE]
#define BS(i, j) Bs[j + i * BLOCK_SIZE]

__kernel void
matrixMul(__global float* A, __global float* B, __global float* C, 
   __local float* As, __local float* Bs, int uiWA, int uiWB)
{
    int bx = get_group_id(0);
    int by = get_group_id(1);
    int tx = get_local_id(0);
    int ty = get_local_id(1);
    int aBegin = uiWA * BLOCK_SIZE * by;
    int aEnd   = aBegin + uiWA - 1;
    int aStep  = BLOCK_SIZE;
    int bBegin = BLOCK_SIZE * bx;
    int bStep  = BLOCK_SIZE * uiWB;
    float Csub = 0.0f;
    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
        AS(ty, tx) = A[a + uiWA * ty + tx];
        BS(ty, tx) = B[b + uiWB * ty + tx];
        barrier(CLK_LOCAL_MEM_FENCE); 
        #pragma unroll
        for (int k = 0; k < BLOCK_SIZE; ++k)
        Csub += AS(ty, k) * BS(k, tx);
            barrier(CLK_LOCAL_MEM_FENCE);
        }
        C[get_global_id(1) * get_global_size(0) + get_global_id(0)] = Csub; 
    }

尺寸为：

#define BLOCK_SIZE 16

size_t localWorkSize[] = {BLOCK_SIZE, BLOCK_SIZE};
size_t globalWorkSize[] = {shrRoundUp(BLOCK_SIZE, N), shrRoundUp(BLOCK_SIZE, M)};

size_t shrRoundUp(int group_size, int global_size) 
{
    int r = global_size % group_size;
    if(r == 0) 
    {
        return global_size;
    } else 
    {
        return global_size + group_size - r;
    }
}

代码取自Nvidia OpenCL矩阵乘法示例。我的 GPU 是：Intel(R) HD Graphics 4600。

谢谢！

Answer 1

我认为你做不到：

cl_mem A_sub = (cl_mem)((float *)A+x*K);

因为cl_mem是OpenCL中的对象，它其实是一个复杂的数据结构，而不仅仅是一个数据指针。它维护诸如指向实际内存的数据指针、对对象的引用、内存属性等信息。不同的运行时代甚至可能有不同的 cl_mem 对象实现。这就是您收到 CL_INVALID_MEM_OBJECT 错误消息的原因。

要获取子矩阵所需的数据，您可以执行以下操作之一：

定义两个新的cl_mem对象，并使用单独的内核来做复制工作。
使用clEnqueueCopyBuffer函数复制主机端的数据代码域。
使用CL_MEM_ALLOC_HOST_PTR内存缓冲，然后使用 clEnqueueMapBuffer 将 GPU 内存映射到主机内存指针，并且然后使用映射的主机内存修改内存内容指针，完成后，取消映射到 return GPU 内存的指针到设备内存域。

OpenCL如何改变cl_mem的内存地址？

OpenCL how to change the memory address of cl_mem?

opencl

matrix-multiplication