OpenCL 矩阵乘法速度

Question

我编写了一个小型 OpenCL 应用程序来计算两个矩阵的乘积。现在我注意到，如果矩阵的大小超过 8192 x 8192，性能会显着下降（16384 x 16384 的计算速度要慢 80 倍），甚至串行实现也要快 5 倍以上。这是主机代码：

/*Make some includes and definitions here*/
#include "stdafx.h"
#include <CL/cl.hpp>

#include <vector>
#include <iostream>

#include "util.hpp" // utility library

#define __CL_ENABLE_EXCEPTIONS
#define ROWS (16384)    // ROWS of vectors a, b, and c
#define COLUMNS (16384)

/*~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~*/
#include "metrics.h"

/*Start main()*/

int main(void)
{
    int A;

    // Fill vectors X and Y with random float values

    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;;
        }
    }
    float* h_y = new float[ROWS*COLUMNS];
    for (int i = 0; i < ROWS; ++i){
        for (int j = 0; j < COLUMNS; ++j){
            h_y[j + i*COLUMNS] = rand() / (float)RAND_MAX;;
        }
    }
    float* h_s = new float[ROWS*COLUMNS];
    for (int i = 0; i < ROWS; ++i){
        for (int j = 0; j < COLUMNS; ++j){
            h_s[j + i*COLUMNS] = 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);

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

    StartCounter();
    //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);

    cl::NDRange global(ROWS*COLUMNS);
    queue.enqueueNDRangeKernel(kernel_add, cl::NullRange, global, cl::NullRange);
    queue.finish();

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

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



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


            std::cout << "" << h_x[i] << "\t ";

    }

    std::cout << "\n\nMatrix #2: \n";
    for (int i = 0; i<ROWS*COLUMNS; i++){


            std::cout << "" << h_y[i] << "\t ";

    }

    std::cout << "\n\nResult: \n";
    for (int i = 0; i<ROWS*COLUMNS; i++){


            std::cout << "" << h_s[i] << "\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)];

}

能解释一下原因吗？我知道如果我执行一些算法优化我可以获得更好的性能，但我试图弄清楚这是否是 "naive" 实现的阈值，或者我做错了什么（不正确的分配分组工作）。

编辑： 因为我在评论中被要求，所以我运行内核的 GPU 是 AMD R9 270/2GB RAM。 CPU 是 i7-4771，系统有 8GB 内存。

Answer 1

写一个关于 "how to do more calculations per thread" 的答案，因为评论中不存在代码格式，并且还涉及一些内存使用...

因此，大多数 OpenCL 实现需要运行每个线程（以及正确的线程数）多条指令才能实现高效性能。但就像我在评论中所说的那样，这在很大程度上取决于处理单元的实际架构（GPU，CPU，或者用独角兽毛编织而成的支持 OpenCL 的魔法单元，无论它是什么）——每个 GPU 制造商， CPUs 和独角兽编织者对如何制作一个非常高效的单元有自己的想法，并且随着时间的流逝，他们也都倾向于改变主意...;)

要在一个线程中完成更多工作，您可以简单地执行以下操作：

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

   for(i = 0; i < NUM_PER_THREAD; i++)
   {
      size_t index = get_global_id(0)*NUM_PER_THREAD + i;
      S[index] = X[index] * Y[index];
   }
}

[这将完成 1 x 16 个块。尝试做 16 x 16 或类似的东西会更有趣，但如果你知道矩阵的大小（宽度）就可以做到]

关于内存：如果所有数据都适合显存，则具有专用本地内存的 GPU（换句话说大多数显卡）的工作速度会快得多。访问 "main" 内存涉及两种方法之一：

当 GPU 通过 PCI-express 总线 [或使用任何基础设施] 读取时，每个缓存行的访问时间很长 - 这可能比 "local" 内存慢 100 或 1000 倍。并且 GPU 还（很可能）必须询问 CPU 内存内容是否在缓存中，如果是，则进一步等待 CPU 将数据复制到主内存...
"page in/out"那里GPU停止，向CPU发送中断， CPU 找到了一些合适的块 [这里的块是 "some amount of memory most likely around 4K or multiple thereof"] 从 GPU 到 "remove" 内存的技术术语内存，并将其复制到主存，然后复制到需要其他内存块到 GPU 内存 - 类似于 OS 交换内存 to/from 硬盘。如果运气不好，GPU 还必须进行一些有趣的缓存或 TLB 刷新，以确保使用正确的数据。

请注意，我仍然（在过去一个小时左右）对 AMD/ATI GPU 的工作方式或其 OpenCL 驱动程序的工作方式没有任何特别的了解。以上内容综合了 guessing/knowing GPU 的一般工作方式、对 OpenCL 的一般工作方式的理解，以及使用 float.[=12 计算存储三个不同的 16K x 16K 数组所需的内存。 =]

OpenCL 矩阵乘法速度

OpenCL Matrix Multiplication Speed

c++

gpu

opencl