【问题标题】:My opencl test does not run much faster than CPU我的 opencl 测试运行速度并不比 CPU 快多少
【发布时间】:2017-07-10 15:30:22
【问题描述】:

我正在尝试测量 GPU 的执行时间并将其与 CPU 进行比较。 我编写了一个 simple_add 函数来添加一个短整型向量的所有元素。 内核代码是:

global const int * A, global const uint * B, global int* C)
    {
        ///------------------------------------------------
        /// Add 16 bits of each
        int AA=A[get_global_id(0)];
        int BB=B[get_global_id(0)];
        int AH=0xFFFF0000 & AA;
        int AL=0x0000FFFF & AA;
        int BH=0xFFFF0000 & BB;
        int BL=0x0000FFFF & BB;
        int CL=(AL+BL)&0x0000FFFF;
        int CH=(AH+BH)&0xFFFF0000;      
        C[get_global_id(0)]=CH|CL;               
     }

我为此函数编写了另一个 CPU 版本,并在执行 100 次后测量了它们的执行时间

clock_t before_GPU = clock();
for(int i=0;i<100;i++)
{
  queue.enqueueNDRangeKernel(kernel_add,1,
  cl::NDRange((size_t)(NumberOfAllElements/4)),cl::NDRange(64));
  queue.finish();
 }
 clock_t after_GPU = clock();


 clock_t before_CPU = clock();
 for(int i=0;i<100;i++)
     AddImagesCPU(A,B,C);
  clock_t after_CPU = clock();

调用整个测量函数10次,结果如下:

        CPU time: 1359
        GPU time: 1372
        ----------------
        CPU time: 1336
        GPU time: 1269
        ----------------
        CPU time: 1436
        GPU time: 1255
        ----------------
        CPU time: 1304
        GPU time: 1266
        ----------------
        CPU time: 1305
        GPU time: 1252
        ----------------
        CPU time: 1313
        GPU time: 1255
        ----------------
        CPU time: 1313
        GPU time: 1253
        ----------------
        CPU time: 1384
        GPU time: 1254
        ----------------
        CPU time: 1300
        GPU time: 1254
        ----------------
        CPU time: 1322
        GPU time: 1254
        ----------------

问题是我真的希望 GPU 比 CPU 快得多,但事实并非如此。我不明白为什么我的 GPU 速度并不比 CPU 高很多。我的代码有问题吗?? 这是我的 GPU 属性:

        -----------------------------------------------------
        ------------- Selected Platform Properties-------------:
        NAME:   AMD Accelerated Parallel Processing
        EXTENSION:      cl_khr_icd cl_amd_event_callback cl_amd_offline_devices cl_khr_d3d10_sharing
        VENDOR:         Advanced Micro Devices, Inc.
        VERSION:        OpenCL 1.2 AMD-APP (937.2)
        PROFILE:        FULL_PROFILE
        -----------------------------------------------------
        ------------- Selected Device Properties-------------:
        NAME :  ATI RV730
        TYPE :  4
        VENDOR :        Advanced Micro Devices, Inc.
        PROFILE :       FULL_PROFILE
        VERSION :       OpenCL 1.0 AMD-APP (937.2)
        EXTENSIONS :    cl_khr_gl_sharing cl_amd_device_attribute_query cl_khr_d3d10_sharing
        MAX_COMPUTE_UNITS :     8
        MAX_WORK_GROUP_SIZE :   128
        OPENCL_C_VERSION :      OpenCL C 1.0
        DRIVER_VERSION:         CAL 1.4.1734
        ==========================================================

只是为了比较这是我的 CPU 规格:

        ------------- CPU Properties-------------:
        NAME :          Intel(R) Core(TM) i3-2100 CPU @ 3.10GHz
        TYPE :  2
        VENDOR :        GenuineIntel
        PROFILE :       FULL_PROFILE
        VERSION :       OpenCL 1.2 AMD-APP (937.2)
        MAX_COMPUTE_UNITS :     4
        MAX_WORK_GROUP_SIZE :   1024
        OPENCL_C_VERSION :      OpenCL C 1.2
        DRIVER_VERSION:         2.0 (sse2,avx)
        ==========================================================

我还使用 QueryPerformanceCounter 测量了挂钟时间,结果如下:

            CPU time: 1304449.6  micro-sec
            GPU time: 1401740.82  micro-sec
            ----------------------
            CPU time: 1620076.55  micro-sec
            GPU time: 1310317.64  micro-sec
            ----------------------
            CPU time: 1468520.44  micro-sec
            GPU time: 1317153.63  micro-sec
            ----------------------
            CPU time: 1304367.29  micro-sec
            GPU time: 1251865.14  micro-sec
            ----------------------
            CPU time: 1301589.17  micro-sec
            GPU time: 1252889.4  micro-sec
            ----------------------
            CPU time: 1294750.21  micro-sec
            GPU time: 1257017.41  micro-sec
            ----------------------
            CPU time: 1297506.93  micro-sec
            GPU time: 1252768.9  micro-sec
            ----------------------
            CPU time: 1293511.29  micro-sec
            GPU time: 1252019.88  micro-sec
            ----------------------
            CPU time: 1320753.54  micro-sec
            GPU time: 1248895.73  micro-sec
            ----------------------
            CPU time: 1296486.95  micro-sec
            GPU time: 1255207.91  micro-sec
            ----------------------

我再次尝试了执行时间的 opencl 分析。

            queue.enqueueNDRangeKernel(kernel_add,1,
                                    cl::NDRange((size_t)(NumberOfAllElements/4)),
                                    cl::NDRange(64),NULL,&ev);
            ev.wait();
            queue.finish();
            time_start=ev.getProfilingInfo<CL_PROFILING_COMMAND_START>();
            time_end=ev.getProfilingInfo<CL_PROFILING_COMMAND_END>();

一次执行的结果大致相同:

            CPU time: 13335.1815  micro-sec
            GPU time: 11865.111  micro-sec
            ----------------------
            CPU time: 13884.0235  micro-sec
            GPU time: 11663.889  micro-sec
            ----------------------
            CPU time: 19724.7296  micro-sec
            GPU time: 14548.222  micro-sec
            ----------------------
            CPU time: 19945.3199  micro-sec
            GPU time: 15331.111  micro-sec
            ----------------------
            CPU time: 17973.5055  micro-sec
            GPU time: 11641.444  micro-sec
            ----------------------
            CPU time: 12652.6683  micro-sec
            GPU time: 11632  micro-sec
            ----------------------
            CPU time: 18875.292  micro-sec
            GPU time: 14783.111  micro-sec
            ----------------------
            CPU time: 32782.033  micro-sec
            GPU time: 11650.444  micro-sec
            ----------------------
            CPU time: 20462.2257  micro-sec
            GPU time: 11647.778  micro-sec
            ----------------------
            CPU time: 14529.6618  micro-sec
            GPU time: 11860.112  micro-sec

【问题讨论】:

  • clock() 测量 CPU 时间而不是挂钟时间。它不会计入 GPU 运行时间。您测量的时间可能由 OpenCL API 调用占用。请尝试 C 中的 clock_gettime() 或 C++ 中的 std::chrono::steady_clock。您没有提到“CPU时间”的单位。如果是clock()函数的原始输出(必须除以CLOCKS_PER_SEC才能得到秒数),1200真的很短。
  • 请参阅this answer 了解 OpenCL 内核时间测量。
  • 由于我是在比较两个执行时间,我认为使用 CPU 时间或挂钟时间并不重要。但是,我尝试以微秒为单位测量挂钟时间并添加此测量值。
  • 内核代码是内存有限的,我怀疑你能优化它。 OpenCL 不适用于此类工作负载。如果此操作是其他数学的前/后阶段,您应该在内核中编写该数学,而不仅仅是位混合步骤。
  • 值得注意的是,GPU 架构通常针对大量浮点运算进行优化,而很少关注整数运算。当 GPU 足够老且工作负载涉及大量整数运算时,我发现在 CPU 上完成奇异工作负载的速度比在 GPU 上快得多。

标签: c++ parallel-processing opencl gpu


【解决方案1】:

我做了一些额外的测试,发现 GPU 已针对浮点运算进行了优化。 我将测试代码更改如下:

void kernel simple_add(global const int * A, global const uint * B, global int* C)
    {
        ///------------------------------------------------
        /// Add 16 bits of each
        int AA=A[get_global_id(0)];
        int BB=B[get_global_id(0)];
        float AH=0xFFFF0000 & AA;
        float AL=0x0000FFFF & AA;
        float BH=0xFFFF0000 & BB;
        float BL=0x0000FFFF & BB;
        int CL=(int)(AL*cos(AL)+BL*sin(BL))&0x0000FFFF;
        int CH=(int)(AH*cos(AH)+BH*sin(BL))&0xFFFF0000;
           C[get_global_id(0)]=CH|CL;               
     }

得到了我预期的结果(快了大约 10 倍):

                CPU time:      741046.665  micro-sec
                GPU time:       54618.889  micro-sec
                ----------------------------------------------------
                CPU time:      741788.112  micro-sec
                GPU time:       54875.666  micro-sec
                ----------------------------------------------------
                CPU time:      739975.979  micro-sec
                GPU time:       54560.445  micro-sec
                ----------------------------------------------------
                CPU time:      755848.937  micro-sec
                GPU time:       54582.111  micro-sec
                ----------------------------------------------------
                CPU time:      724100.716  micro-sec
                GPU time:       56893.445  micro-sec
                ----------------------------------------------------
                CPU time:      744476.351  micro-sec
                GPU time:       54596.778  micro-sec
                ----------------------------------------------------
                CPU time:      727787.538  micro-sec
                GPU time:       54602.445  micro-sec
                ----------------------------------------------------
                CPU time:      731132.939  micro-sec
                GPU time:       54710.000  micro-sec
                ----------------------------------------------------
                CPU time:      727899.150  micro-sec
                GPU time:       54583.444  micro-sec
                ----------------------------------------------------
                CPU time:      727089.880  micro-sec
                GPU time:       54594.778  micro-sec
                ----------------------------------------------------

对于更重的浮点运算,如下所示:

        void kernel simple_add(global const int * A, global const uint * B, global int* C)
            {
                ///------------------------------------------------
                /// Add 16 bits of each
                int AA=A[get_global_id(0)];
                int BB=B[get_global_id(0)];
                float AH=0xFFFF0000 & AA;
                float AL=0x0000FFFF & AA;
                float BH=0xFFFF0000 & BB;
                float BL=0x0000FFFF & BB;
                int CL=(int)(AL*(cos(AL)+sin(2*AL)+cos(3*AL)+sin(4*AL)+cos(5*AL)+sin(6*AL))+
                        BL*(cos(BL)+sin(2*BL)+cos(3*BL)+sin(4*BL)+cos(5*BL)+sin(6*BL)))&0x0000FFFF;
                int CH=(int)(AH*(cos(AH)+sin(2*AH)+cos(3*AH)+sin(4*AH)+cos(5*AH)+sin(6*AH))+
                        BH*(cos(BH)+sin(2*BH)+cos(3*BH)+sin(4*BH)+cos(5*BH)+sin(6*BH)))&0xFFFF0000;
                        C[get_global_id(0)]=CH|CL;

             }

结果差不多:

                CPU time:     3905725.933  micro-sec
                GPU time:      354543.111  micro-sec
                -----------------------------------------
                CPU time:     3698211.308  micro-sec
                GPU time:      354850.333  micro-sec
                -----------------------------------------
                CPU time:     3696179.243  micro-sec
                GPU time:      354302.667  micro-sec
                -----------------------------------------
                CPU time:     3692988.914  micro-sec
                GPU time:      354764.111  micro-sec
                -----------------------------------------
                CPU time:     3699645.146  micro-sec
                GPU time:      354287.666  micro-sec
                -----------------------------------------
                CPU time:     3681591.964  micro-sec
                GPU time:      357071.889  micro-sec
                -----------------------------------------
                CPU time:     3744179.707  micro-sec
                GPU time:      354249.444  micro-sec
                -----------------------------------------
                CPU time:     3704143.214  micro-sec
                GPU time:      354934.111  micro-sec
                -----------------------------------------
                CPU time:     3667518.628  micro-sec
                GPU time:      354809.222  micro-sec
                -----------------------------------------
                CPU time:     3714312.759  micro-sec
                GPU time:      354883.888  micro-sec
                -----------------------------------------

【讨论】:

    【解决方案2】:

    ATI RV730 具有 VLIW 结构,因此最好尝试具有 1/4 总线程数(即 NumberOfAllElements/16)的 uint4int4 向量类型。这也有助于更快地从内存中加载每个工作项。

    与内存操作相比,内核也没有太多的计算。将缓冲区映射到 RAM 会有更好的性能。不要复制数组,使用 map/unmap enqueue 命令将它们映射到内存。

    如果还是不快,你可​​以同时使用 gpu 和 cpu 来处理前半部分和后半部分的工作,以在 %50 时间内完成。

    也不要将 clFinish 放入循环中。把它放在循环结束之后。这样,它将更快地加入队列,并且已经按顺序执行,因此在完成第一个项目之前它不会启动其他项目。我想这是有序队列,并且在每个入队之后添加 cfinish 是额外的开销。在最新内核之后只需要一个 clfinish 就足够了。


    ATI RV730:64 个 VLIW 单元,每个单元至少有 4 个流核心。 750 兆赫。

    i3-2100:2 个内核(仅用于防冒泡的线程),每个内核都具有能够同时流式传输 8 个操作的 AVX。因此,这可以进行 16 次操作。超过 3 GHz。

    流操作与频率的简单乘法:

    ATI RV730 = 192 个单位(更多具有乘加功能,每个 vliw 的第 5 个元素)

    i3-2100 = 48 个单位

    所以 gpu 应该至少快 4 倍(使用 int4、uint4)。这适用于简单的 ALU 和 FPU 操作,例如按位操作和乘法。诸如超验性能之类的特殊功能可能会有所不同,因为它们仅在每个 vliw 的第 5 个单元上运行。

    【讨论】:

    • 我在测量中没有考虑数据传输时间。我在复制后胡思乱想,gpu 必须执行更快的比较内存映射(我还不确定)
    • 那么它是 vliw 微架构需要 4 宽向量而不是您使用的标量。如果有 1M 线程,现在只有 256k 线程可以使用 int4 更快地完成工作
    • 从循环中删除clfinish,放在循环之后
    • 上次使用没有循环的 opencl 分析(一次)。结果没有改变(我将它们添加到问题中)。
    • @Afshin 运行一次意味着它还没有优化。多次运行后会好转。入队 10 次,最后 clfinish 一次
    猜你喜欢
    • 2015-09-23
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 2014-06-15
    • 2011-05-15
    • 2016-10-23
    • 2017-03-20
    相关资源
    最近更新 更多