OpenCL - 全局内存读取性能优于本地

Posted 2023-02-16

【中文标题】OpenCL - 全局内存读取性能优于本地

【英文标题】：OpenCL - Global Memory reads preforming better than local

【发布时间】：2012-09-17 22:53:07

【问题描述】：

我有一个在 NVidia GTX 680 上运行的内核，当从使用全局内存切换到本地内存时，它会增加执行时间。

我的内核是有限元光线追踪器的一部分，现在在处理之前将每个元素加载到本地内存中。每个元素的数据存储在结构 fastTriangle 中，其定义如下：

typedef struct fastTriangle 
    float cx, cy, cz, cw;
    float nx, ny, nz, nd;
    float ux, uy, uz, ud;
    float vx, vy, vz, vd;
 fastTriangle;

我将这些对象的数组传递给内核，编写如下（为简洁起见，我删除了不相关的代码：

__kernel void testGPU(int n_samples, const int n_objects, global const fastTriangle *objects, __local int *x_res, __global int *hits) 
    // Get gid, lid, and lsize

    // Set up random number generator and thread variables

    // Local storage for the two triangles being processed
    __local fastTriangle triangles[2]; 

    for(int i = 0; i < n_objects; i++)     // Fire ray from each object
        event_t evt = async_work_group_copy((local float*)&triangles[0], (global float*)&objects[i],sizeof(fastTriangle)/sizeof(float),0);

        //Initialise local memory x_res to 0's

        barrier(CLK_LOCAL_MEM_FENCE);
        wait_group_events(1, &evt);      


        Vector wsNormal =  triangles[0].cw*triangles[0].nx, triangles[0].cw*triangles[0].ny, triangles[0].cw*triangles[0].nz;

        for(int j = 0; j < n_samples; j+= 4) 
            // generate a float4 of random numbers here (rands

            for(int v = 0; v < 4; v++)     // For each ray in ray packet
                //load the first object to be intesected
                evt = async_work_group_copy((local float*)&triangles[1], (global float*)&objects[0],sizeof(fastTriangle)/sizeof(float),0);

                // Some initialising code and calculate ray here
                // Should have ray fully specified at this point;


                for(int w = 0; w < n_objects; w++)         // Check for intersection against each ray

                    wait_group_events(1, &evt);

                    // Check for intersection against object w


                    float det = wsDir.x*triangles[1].nx + wsDir.y*triangles[1].ny + wsDir.z*triangles[1].nz;
                    float dett = triangles[1].nd - (triangles[0].cx*triangles[1].nx + triangles[0].cy*triangles[1].ny + triangles[0].cz*triangles[1].nz);


                    float detpx = det*triangles[0].cx + dett*wsDir.x;
                    float detpy = det*triangles[0].cy + dett*wsDir.y;
                    float detpz = det*triangles[0].cz + dett*wsDir.z;


                    float detu = detpx*triangles[1].ux + detpy*triangles[1].uy + detpz*triangles[1].uz + det*triangles[1].ud;
                    float detv = detpx*triangles[1].vx + detpy*triangles[1].vy + detpz*triangles[1].vz + det*triangles[1].vd;


                    // Interleaving the copy of the next triangle
                    evt = async_work_group_copy((local float*)&triangles[1], (global float*)&objects[w+1],sizeof(fastTriangle)/sizeof(float),0);

                    // Complete intersection calculations

                 // end for each object intersected

                if(objectNo != -1) atomic_inc(&x_res[objectNo]);
             // end for sub rays
         // end for each ray
        barrier(CLK_LOCAL_MEM_FENCE);

        // Add all the local x_res to global array hits


        barrier(CLK_GLOBAL_MEM_FENCE);
     // end for each object

当我第一次编写这个内核时，我没有在本地内存中缓冲每个对象，而是从全局内存中访问它，即我将使用 objects[i].cx 而不是 triangles[0].cx

在开始优化时，我切换到使用上面列出的本地内存，但随后观察到执行运行时间增加了大约 25%。

为什么使用本地内存缓冲对象而不是直接在全局内存中访问它们时性能会更差？

【参考方案1】：

本地内存是否可以帮助您更快地运行，这实际上取决于您的程序。使用本地内存有两点需要考虑：

在将数据从全局复制到本地以及再次从本地复制到全局时，您需要进行额外的计算。

我看到你有 3 次“障碍（...）”，这些障碍是性能杀手。所有 OpenCL 任务都必须等待所有其他任务。这样会阻碍并行性，任务不再独立运行。

当您在计算中多次读取数据时，本地内存非常有用。但是，与复制和同步相比，快速读取和写入需要为您带来更多的性能提升。

【讨论】：

对我来说很清楚。他正在从全局复制到本地，使用一次并删除数据。 Fr 这个目的最好是 1 次访问。