OpenCL矩阵乘法速度

2022-03-27 00:00:00 gpu opencl c++

我编写了一个计算两个矩阵乘积的小型OpenCL应用程序。现在我注意到，如果矩阵的大小超过8192x8192，则性能会显著下降(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!
";
        exit(1);
    }

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

    // 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!
";
        exit(1);
    }

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

    // 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) << "
";
        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 
";

    //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 << "
Matrix #1: 
";
    for (int i = 0; i<ROWS*COLUMNS; i++){


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

    }

    std::cout << "

Matrix #2: 
";
    for (int i = 0; i<ROWS*COLUMNS; i++){


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

    }

    std::cout << "

Result: 
";
    for (int i = 0; i<ROWS*COLUMNS; i++){


            std::cout << "" << h_s[i] << "	 ";

    }*/
    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)];

}

您能给我解释一下原因吗？我知道如果我执行一些算法优化，我可以获得更好的性能，但是我想知道这是"幼稚"实现的门槛，还是我做错了什么(不正确地将工作分配给组)。

edit：因为我在注释中被要求使用GPU，所以我运行内核的GPU是AMDR9 270/2 GB RAM。CPU为i7-4771，系统具有8 GB RAM。

解决方案

写一个关于"如何在每个线程进行更多计算"的答案，因为注释中不存在代码格式化，并且还涉及一些关于内存使用的内容.

因此，大多数OpenCL实现需要在每个线程(以及适当数量的线程)上运行几条以上的指令才能获得高效的性能。但就像我在评论中所说的，这高度依赖于处理单元的实际架构(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];
   }
}

[这将执行1x16个块。尝试做16x16或类似的事情会更有趣，但如果您知道矩阵的大小(宽度)也可以做到]

关于内存：拥有专用本地内存的GPU(换句话说，大多数显卡)的工作速度会快得多，如果所有数据都能放在图形内存中的话。访问"主"内存涉及以下两种方法之一：

  1. 当GPU通过PCI-Express总线[或使用的任何基础设施]读取时，每个高速缓存线的访问时间都很长--这可能比本地内存慢100倍或1000倍。而GPU还(很可能)必须询问CPU内存内容是否在缓存中，如果在缓存中，则进一步等待CPU将数据复制到主存中……
  2. GPU停止的"Page In/Out"，向CPU发送中断， CPU会从GPU中找到一些合适的内存块[这里的Lump是一个术语，指的是很可能是4K左右或其倍数的某个内存量] 内存，并将其复制到主内存，然后复制到 GPU内存所需的其他内存块-类似于操作系统在硬盘之间交换内存时的情况。如果您运气不佳，GPU还必须执行一些有趣的缓存或TLB刷新，以确保正在使用正确的数据。

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

相关文章