c++ - 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)];

}

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

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

Answer 1

写一个关于“如何为每个线程做更多计算”的答案，因为注释中不存在代码格式，并且还涵盖了一些内存使用...

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

要在一个线程中做更多的工作，你可以简单地做:

#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 个 block 。尝试做 16 x 16 或类似的东西会更有趣，但如果你知道矩阵的大小(宽度)就可以做到]

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

1. 当 GPU 通过 PCI-express 总线 [或任何使用的基础设施] 读取时，每个缓存行的访问时间很长 - 这可能比“本地”内存慢 100 或 1000 倍。并且 GPU 还(很可能)必须询问 CPU 内存内容是否在缓存中，如果是，则进一步等待 CPU 将数据复制到主内存...
2. GPU 停止的“页面进/出”，向 CPU 发送中断，CPU 找到一些合适的内存块 [在此上下文中的 block 是“最有可能大约 4K 或其倍数的内存量”] 的内存块以从 GPU 中“移除”内存，并将其复制到主存，然后复制到需要其他内存块到 GPU 内存 - 类似于操作系统将内存交换到硬盘或从硬盘交换内存时。如果运气不好，GPU 还必须进行一些有趣的缓存或 TLB 刷新，以确保使用正确的数据。

请注意，我仍然(在过去一小时左右)对 AMD/ATI GPU 的工作方式或其 OpenCL 驱动程序的工作方式没有任何特别的了解。以上是猜测/了解 GPU 的一般工作方式、了解 OpenCL 的一般工作方式以及使用 float 计算存储三个不同的 16K x 16K 阵列所需的内存的混合。 .