OpenCL 矩阵乘法速度

Posted 2023-02-16

【中文标题】OpenCL 矩阵乘法速度

【英文标题】：OpenCL Matrix Multiplication Speed

【发布时间】：2015-07-15 19:35:53

【问题描述】：

我编写了一个计算两个矩阵乘积的小型 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)];

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

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

【问题讨论】：

虽然 OpenCL 非常便携，但它并不总是“性能便携”，因此您可能应该告诉我们您正在运行什么硬件。在大多数架构上，每个内核线程执行一次乘法可能非常慢。我会做 16 x 16、4 x 32 或 256 x 256 之类的东西。“正确”的大小取决于您的硬件，但每个线程中的线程少一点多一点可能会更好。检查您的 OpenCL 实现的文档是否说明了“调整内核”等主题。

因此，如果您有 2GB 的专用图形内存，则不会存储 3 x sizeof(float) x 16384 x 16384，因为这会占用 3GB。这可以解释为什么它很慢 - 我不确定 AMD/ATI OpenCL 驱动程序在这种情况下会做什么，但以一种或另一种方式，它必须至少将一些分配存储在主内存中，或者使用GPU 分页以一次又一次地交换它，或者使用某种总线来访问 CPU 内存空间中的数据。其中任何一个都会让它变得很慢，我预计[我对 ATI/AMD 在他们的产品中所做的事情一无所知]

@MatsPetersson 感谢您的建议。我很想知道这是由于内存问题还是因为每个线程的计算很少。但是我应该怎么做才能改变每个线程的计算呢？

您的主机代码也在计时从设备传回的数据，这可能在您测量的运行时间中占主导地位。为了对这类事情进行基准测试，我通常会在一个循环中多次运行内核并取平均值，这让设备有机会预热。

是的，我已经看到并更正了 - 我也应该在这里更正。对于 8192x8192（约 50 毫秒内核执行时间），整个交易的速度大约快了几毫秒，而对于更大的 16384x16384（内核执行时间大约为 7 秒），整个交易大约快了 0.5 秒。现在，关于您的建议，我将对其进行测试，看看会发生什么。

【参考方案1】：

写一个关于“如何为每个线程进行更多计算”的答案，因为在 cmets 中不存在代码格式化，并且还涵盖了一点内存使用情况......

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

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

#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 将运行得更快。访问“主”内存涉及以下两种方法之一：

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

请注意，我仍然（在最后一个小时左右）对 AMD/ATI GPU 的工作原理或 OpenCL 驱动程序的工作原理没有任何特别的了解。以上是猜测/了解 GPU 一般如何工作、了解 OpenCL 一般如何工作以及使用float 计算存储三个不同的 16K x 16K 数组所需的内存的混合体。

【讨论】：

有趣的是，我没有看到任何性能提升。我已经尝试了许多“结构”，例如不同的全局和局部工作大小、不同的线程数和每个线程的元素计算，但没有任何显着的结果。所以我认为正如你所指出的，可能是内存不足的问题。

是的，几乎可以肯定，“没有足够的 GPU 内存”是问题所在，当然，如果您有 2GB 的 GPU 内存并使用该 GPU 在您的计算机上绘制屏幕，​​它会可能会使用一点 GPU 内存作为屏幕帧缓冲区以及在计算机屏幕上绘制所需的各种其他东西 - 字体、图标位图等