OpenCL 二维数组乘法

Question

我刚刚开始尝试使用 OpenCL。我正在尝试制作一个将两个二维数组相乘的内核。我已经用向量做到了这一点，但是在二维中我只得到第一行的结果。我已经尝试实施一些我发现的解决方案，但它们中的每一个都只与第一行混淆。 执行过程中的图像： http://i.imgur.com/lJqSURV.png

这是我的主机文件：

#include "stdafx.h"
#include <CL/cl.hpp>

#include <vector>
#include <iostream>

#include "util.hpp" // utility library   

#define __CL_ENABLE_EXCEPTIONS
#define ROWS (5)
#define COLUMNS (5)

#include "metrics.h"

/*Start main()*/

int main(void)
{
    int A = 4;
    /*Define the vectors for operands and result*/

    float** h_x = new float*[ROWS];
    float** h_y = new float*[ROWS];
    float** h_s = new float*[ROWS];

    for (int i = 0; i < ROWS; ++i){
        h_x[i] = new float[COLUMNS];
    }

    for (int i = 0; i < ROWS; ++i){
        h_y[i] = new float[COLUMNS];
    }

    for (int i = 0; i < ROWS; ++i){
        h_s[i] = new float[COLUMNS];
    }

    // Fill vectors a and b with random float values

    for (int i = 0; i < ROWS; i++)
    {
        for (int j = 0; j < COLUMNS; j++){
            h_x[i][j] = rand() / (float)RAND_MAX;
            h_y[i][j] = rand() / (float)RAND_MAX;
            h_s[i][j] = 0.0;
        }   
    }

    /*~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~*/

    // Get all platforms (drivers)

    std::vector<cl::Platform> all_platforms;
    cl::Platform::get(&all_platforms);


    if (all_platforms.size() == 0){ // Check for issues
        std::cout << " No platforms found. Check OpenCL installation!\n";
        exit(1);
    }

    cl::Platform default_platform = all_platforms[0];
    std::cout << "Using platform: " << default_platform.getInfo<CL_PLATFORM_NAME>() << "\n";

    // Get default device of the default platform

    std::vector<cl::Device> all_devices;
    default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);

    if (all_devices.size() == 0){ // Check for issues
        std::cout << " No devices found. Check OpenCL installation!\n";
        exit(1);
    }

    cl::Device default_device = all_devices[0];
    std::cout << "Using device: " << default_device.getInfo<CL_DEVICE_NAME>() << "\n";

    // Create an OpenCL context

    cl::Context context({ default_device });

    cl::Program program(context, util::loadProgram("saxy_kernel.cl"), true);

    if (program.build({ default_device }) != CL_SUCCESS){
        std::cout << " Error building: " << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device) << "\n";
        getchar();
        exit(1);
    }

    // create buffers on the device
    cl::Buffer buffer_X(context, CL_MEM_READ_WRITE, sizeof(float)* ROWS*COLUMNS);
    cl::Buffer buffer_Y(context, CL_MEM_READ_WRITE, sizeof(float)* ROWS*COLUMNS);
    cl::Buffer buffer_S(context, CL_MEM_READ_WRITE, sizeof(float)* ROWS*COLUMNS);
    cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int));

    //create queue to which we will push commands for the device.
    cl::CommandQueue queue(context, default_device);


    StartCounter();
    //write arrays A and B to the device
    queue.enqueueWriteBuffer(buffer_X, CL_TRUE, 0, sizeof(float)* ROWS*COLUMNS, &h_x[0][0]);
    queue.enqueueWriteBuffer(buffer_Y, CL_TRUE, 0, sizeof(float)* ROWS*COLUMNS, &h_y[0][0]);
    queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int), &A);

    //run the kernel
    cl::Kernel kernel_add = cl::Kernel(program, "simple_add");
    kernel_add.setArg(0, buffer_X);
    kernel_add.setArg(1, buffer_Y);
    kernel_add.setArg(2, buffer_S);
    kernel_add.setArg(3, buffer_A);

    queue.enqueueNDRangeKernel(kernel_add, cl::NullRange, cl::NDRange(5,5), cl::NullRange);
    queue.finish();

    //read result C from the device to array C
    queue.enqueueReadBuffer(buffer_S, CL_TRUE, 0, sizeof(float)* ROWS * COLUMNS, &h_s[0][0]);

    std::cout << "Kernel execution time: " << GetCounter() << "ms \n";

    /*Print vectors*/
    std::cout << "\nMatrix #1: \n";
    for (int i = 0; i<ROWS; i++){
        std::cout << "\n";
        for (int j = 0; j<COLUMNS; j++){
            std::cout << "" << h_x[i][j] << "\t ";
        }
    }

    std::cout << "\n\nMatrix #2: \n";
    for (int i = 0; i<ROWS; i++){
        std::cout << "\n";
        for (int j = 0; j<COLUMNS; j++){
            std::cout << "" << h_y[i][j] << "\t ";
        }
    }

    std::cout << "\n\nResult: \n";
    for (int i = 0; i<ROWS; i++){
        std::cout << "\n";
        for (int j = 0; j<COLUMNS; j++){
            std::cout << "" << h_s[i][j] << "\t ";
        }
    }
    getchar();
    return 0;
}

这里是内核：

__kernel void kernel simple_add(
   __global float* X, 
   __global float* Y, 
   __global float* S, 
   __global int *A){

   S[get_global_id(0)] = X[get_global_id(0)] * Y[get_global_id(0)];

/* Var defs
   int k;
   int i = get_global_id(0);
   int j = get_global_id(1);
   float tmp;

   if ( (i < 5) && (j < 5))
   {
       tmp = 0.0;
       for(k=0;k<5;k++)
           tmp += X[i*5+k] * Y[k*5+j];
       S[i*5+j] = tmp;
   }*/
}

我确定我做错了什么真的，但我不知道是什么。任何帮助将不胜感激。

Answer 1

您的内核代码很好，就像您创建 OpenCL 缓冲区和启动内核的方式一样。问题在于您的数据在主机上的表示方式，以及您如何将其复制到设备上。

您的 OpenCL 缓冲区是一维数组，这是必需的。但是，您的主机数组是 2D 的，这意味着相邻的行是不连续的（2D 数组是指针数组）。

（最简单的）解决方法是将主机上的存储线性化，以匹配设备的数据布局：

float* h_x = new float[ROWS*COLUMNS];
for (int i = 0; i < ROWS; ++i){
    for (int j = 0; j < COLUMNS; ++j){
      h_x[j + i*COLUMNS] = rand() / (float)RAND_MAX;;
    }
}