OpenCL - 全局內存讀取預製比本地更好

Question

我有一個運行在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 
} 

當我第一次寫這個內核我沒有緩衝在本地內存中的每個對象，而是剛剛訪問它形成全局內存即不是三角形的[0] .CX我會使用對象[I] .CX

當開始優化時，我切換到使用上面列出的本地內存，但觀察到執行運行時間增加了大約25％。

當使用本地內存緩衝對象而不是直接在全局內存中訪問它們時，爲什麼性能會更差？

Answer 1

如果本地內存可以幫助您運行得更快，那它確實取決於您的程序。有使用本地內存時要考慮兩件事情：

1. 複製從全球到地方，並從地方到全球再度數據時，你有額外的計算。

2. 我看到你有3次「障礙（......）」，這些障礙是性能殺手。所有OpenCL任務都必須等待所有其他人的障礙。這種並行性受到阻礙，任務不再獨立運行。

當您在計算中讀取數據時，本地內存很棒。但快速讀取和寫入需要比複製和同步需要更多的性能增益。