OpenCL 中逐元素矩阵乘法速度

问题描述 投票：0回答：1

我编写了一个小型 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 RAM。

c++ opencl gpu
1个回答
2
投票

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

因此，大多数 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];
   }
}

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

关于内存：如果所有数据都适合图形内存，则具有专用本地内存的 GPU（即大多数显卡）的工作速度会快得多。访问“主”内存涉及两种方法之一：

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

请注意，我仍然（在过去一个小时左右）对 AMD/ATI GPU 的工作原理或他们的 OpenCL 驱动程序的工作原理没有任何特别的了解。以上是猜测/了解 GPU 一般如何工作、了解 OpenCL 一般如何工作以及使用 

float
计算存储 16K x 16K 三个不同数组所需的内存的混合体。

最新问题
© www.soinside.com 2019 - 2024. All rights reserved.