Конфликт банка локальной памяти Mobile OpenCL. Почему использование локальной памяти медленнее, чем использование глобальной памяти в ядре?

Я разрабатываю приложение для распознавания лиц на платформе Android, используя OpenCL. Алгоритм распознавания лиц основан на алгоритме Виолы Джонса. Я попытался сделать код ядра шага каскадной классификации. и я установил classifier data каскадной стадии 1 среди каскадных стадий на local memory(__local), потому что данные классификатора используются для всех рабочих элементов.

Но время профилирования ядра без использования локальной памяти (с использованием глобальной памяти) быстрее, чем с использованием локальной памяти.

отредактировано:

Я загрузил полный код.

с локальной версией памяти

__kernel void CASCADE(__read_only image2d_t input_image, __write_only image2d_t output_image,__constant float* classifierMem,__constant int* idxNumValStageArray, int numTotStage, __constant int* vecSkin){
       int cascadeLocalSize = get_local_size(0);

       __local float localS1F1[42];

       int localIdx = get_local_id(1)*cascadeLocalSize + get_local_id(0);
       if(localIdx<42)
       {
           int stage1Idx = localIdx + idxNumValStageArray[0]+4;
           localS1F1[localIdx] = classifierMem[stage1Idx];
       }
       barrier(CLK_LOCAL_MEM_FENCE);


       float resizeFactor = 1.0;
       int2 im_dim = get_image_dim(input_image);
       unsigned int srcWidth = im_dim.x*(float)resizeFactor;
       unsigned int srcHeight = im_dim.y*(float)resizeFactor;

       int gx = get_global_id(0);
       int gy = get_global_id(1);

       int skinX=0;
       int skinY=0;
       int coordi=vecSkin[512*gy+gx];
       skinX = coordi%im_dim.x;
       skinY = coordi/im_dim.x;

       if( skinX >= 10 && skinY >= 10 )
       {
             skinX -= 10;
             skinY -= 10;
       }      

       int type = gx%3;

       unsigned int windowWidth = classifierMem[0];
       unsigned int windowHeight = classifierMem[1]; 


       unsigned int stageIndex;
       float stageThres;
       float numFeatures;
       unsigned int featureIndex;
       float featureValue;

       if(skinX<srcWidth-windowWidth-1 && skinY<srcHeight-windowHeight-1){
             bool stagePass = true;
             unsigned int index = 0;
             for(unsigned int i=numTotStage; i>0;i--){
                    if(stagePass){
                           if(index == 0){
                                 stageIndex = idxNumValStageArray[0];                                 
                                 stageThres = classifierMem[stageIndex+2];
                                 numFeatures = classifierMem[stageIndex+3];
                                 featureIndex = 0;
                                 featureValue = 0.0;                           
                           }
                           else{
                                 stageIndex = idxNumValStageArray[index];
                                 stageThres = classifierMem[stageIndex+2];
                                 numFeatures = classifierMem[stageIndex+3];
                                 featureIndex = stageIndex+4;
                                 featureValue = 0.0;
                           }
                           float featureThres;
                           float succVal;
                           float failVal;
                           unsigned int numRegions;
                           float regionValue;


                           if(type ==0 && index==0)
                           {
                                 for(unsigned int j=numFeatures; j>0;j--){
                                        if(stagePass){
                                               featureThres=localS1F1[featureIndex++]*(windowWidth*windowHeight);
                                              succVal=localS1F1[featureIndex++];
                                              failVal=localS1F1[featureIndex++];
                                              numRegions = localS1F1[featureIndex++];
                                              regionValue =0.0;

                                              float4 rectValue;
                                              int4 regionP;                                  
                                              for(unsigned int k=numRegions; k>0;k--){

                                                     regionP.x=(int)(localS1F1[featureIndex])+skinX;
                                                     regionP.y=(int)(localS1F1[featureIndex+1])+skinY;
                                                     regionP.z=(int)(localS1F1[featureIndex+2])+regionP.x;
                                                     regionP.w=(int)(localS1F1[featureIndex+3])+regionP.y;

                                                     rectValue.x = read_imagef(input_image, sampler, regionP.xy).x;
                                                     rectValue.y = read_imagef(input_image, sampler, regionP.zy).x;
                                                     rectValue.z = read_imagef(input_image, sampler, regionP.xw).x;
                                                     rectValue.w = read_imagef(input_image, sampler, regionP.zw).x;

                                                     regionValue += dot(rectValue, (float4)(1.0f, -1.0f, -1.0f, 1.0f))*localS1F1[featureIndex+4];

                                                     featureIndex+=5;
                                              }
                                              featureValue += (regionValue < featureThres)?failVal:succVal;                              
                                              if(j*2 == (unsigned int)numFeatures && featureValue*2 < stageThres) stagePass =false;

                                        }// end of if(stagePass) 
                                 }// end of for(unsigned int j=numFeatures; j>0;j--)

                                  index++;
                                 if(featureValue < stageThres)    stagePass =false;
                                 else if(index==numTotStage) write_imagef(output_image, (int2)(skinX, skinY), (0.1));

                           }

                           else if(type ==1 && index ==0)
                           {
                                 featureIndex +=14;
                                 for(unsigned int j=numFeatures; j>0;j--){
                                        if(stagePass){
                                              if(j==1)
                                                     featureIndex -= 42;

                                               featureThres=localS1F1[featureIndex++]*(windowWidth*windowHeight);
                                              succVal=localS1F1[featureIndex++];
                                              failVal=localS1F1[featureIndex++];
                                              numRegions = localS1F1[featureIndex++];
                                              regionValue =0.0;


                                              float4 rectValue;
                                              int4 regionP;                                  
                                              for(unsigned int k=numRegions; k>0;k--){

                                                     regionP.x=(int)(localS1F1[featureIndex])+skinX;
                                                     regionP.y=(int)(localS1F1[featureIndex+1])+skinY;
                                                     regionP.z=(int)(localS1F1[featureIndex+2])+regionP.x;
                                                     regionP.w=(int)(localS1F1[featureIndex+3])+regionP.y;

                                                     rectValue.x = read_imagef(input_image, sampler, regionP.xy).x;
                                                     rectValue.y = read_imagef(input_image, sampler, regionP.zy).x;
                                                     rectValue.z = read_imagef(input_image, sampler, regionP.xw).x;
                                                     rectValue.w = read_imagef(input_image, sampler, regionP.zw).x;

                                                     regionValue += dot(rectValue, (float4)(1.0f, -1.0f, -1.0f, 1.0f))*localS1F1[featureIndex+4];

                                                     featureIndex+=5;
                                              }
                                              featureValue += (regionValue < featureThres)?failVal:succVal;                              
                                              if(j*2 == (unsigned int)numFeatures && featureValue*2 < stageThres) stagePass =false;
                                        }
                                 }

                                  index++;
                                 if(featureValue < stageThres)    stagePass =false;
                                 else if(index==numTotStage) write_imagef(output_image, (int2)(skinX, skinY), (0.1));
                           }

                           else if(index == 0)
                           {
                                 featureIndex +=28;
                                 for(unsigned int j=numFeatures; j>0;j--){
                                        if(stagePass){

                                              if(j==2)     featureIndex -= 42;

                                               featureThres=localS1F1[featureIndex++]*(windowWidth*windowHeight);
                                              succVal=localS1F1[featureIndex++];
                                              failVal=localS1F1[featureIndex++];
                                              numRegions = localS1F1[featureIndex++];
                                              regionValue =0.0;

                                              float4 rectValue;
                                              int4 regionP;                                  
                                              for(unsigned int k=numRegions; k>0;k--){

                                                     regionP.x=(int)(localS1F1[featureIndex])+skinX;
                                                     regionP.y=(int)(localS1F1[featureIndex+1])+skinY;
                                                     regionP.z=(int)(localS1F1[featureIndex+2])+regionP.x;
                                                     regionP.w=(int)(localS1F1[featureIndex+3])+regionP.y;

                                                     rectValue.x = read_imagef(input_image, sampler, regionP.xy).x;
                                                     rectValue.y = read_imagef(input_image, sampler, regionP.zy).x;
                                                     rectValue.z = read_imagef(input_image, sampler, regionP.xw).x;
                                                     rectValue.w = read_imagef(input_image, sampler, regionP.zw).x;

                                                     regionValue += dot(rectValue, (float4)(1.0f, -1.0f, -1.0f, 1.0f))*localS1F1[featureIndex+4];

                                                     featureIndex+=5;
                                              }// end of for(unsigned int k=numRegions; k>0;k--)
                                              featureValue += (regionValue < featureThres)?failVal:succVal;                              
                                              if(j*2 == (unsigned int)numFeatures && featureValue*2 < stageThres) stagePass =false;

                                        }// end of if(stagePass)
                                 }//end of for(unsigned int j=numFeatures; j>0;j--)

                                 index++;
                                 if(featureValue < stageThres)    stagePass =false;
                                 else if(index==numTotStage) write_imagef(output_image, (int2)(skinX, skinY), (0.1));
                           }

                           //stage 
                           else{
                                 for(unsigned int j=numFeatures; j>0;j--){
                                        if(stagePass){
                                               featureThres=classifierMem[featureIndex++]*(windowWidth*windowHeight);
                                              succVal=classifierMem[featureIndex++];
                                              failVal=classifierMem[featureIndex++];
                                              numRegions = classifierMem[featureIndex++];
                                              regionValue =0.0;
                                              float4 rectValue;
                                              int4 regionP;                                  
                                              for(unsigned int k=numRegions; k>0;k--){
                                                     regionP.x=(int)(classifierMem[featureIndex])+skinX;
                                                     regionP.y=(int)(classifierMem[featureIndex+1])+skinY;
                                                     regionP.z=(int)(classifierMem[featureIndex+2])+regionP.x;
                                                     regionP.w=(int)(classifierMem[featureIndex+3])+regionP.y;
                                                     rectValue.x = read_imagef(input_image, sampler, regionP.xy).x;
                                                     rectValue.y = read_imagef(input_image, sampler, regionP.zy).x;
                                                     rectValue.z = read_imagef(input_image, sampler, regionP.xw).x;
                                                     rectValue.w = read_imagef(input_image, sampler, regionP.zw).x;
                                                     regionValue += dot(rectValue, (float4)(1.0f, -1.0f, -1.0f, 1.0f))*classifierMem[featureIndex+4]; 
                                                     featureIndex+=5;
                                              }
                                              featureValue += (regionValue < featureThres)?failVal:succVal;                              
                                              if(j*2 == (unsigned int)numFeatures && featureValue*2 < stageThres) stagePass =false;
                                        }
                                 }
                                 index++;
                                 if(featureValue < stageThres)    stagePass =false;
                                 else if(index==numTotStage) write_imagef(output_image, (int2)(skinX, skinY), (0.1));
                           }
                    }
             }      
       }else return;
}

оригинальная версия (без локальной памяти)

__kernel void CASCADE(__read_only image2d_t input_image, __write_only image2d_t output_image,__constant float* classifierMem,__constant int* idxNumValStageArray, int numTotStage, __constant int* vecSkin){
    float resizeFactor = 1.0;

    int2 im_dim = get_image_dim(input_image);

    unsigned int srcWidth = im_dim.x*(float)resizeFactor;
    unsigned int srcHeight = im_dim.y*(float)resizeFactor;

    int gx = get_global_id(0);
    int gy = get_global_id(1);


    int skinX=0;
    int skinY=0;
    int coordi=vecSkin[512*gy+gx];
    skinX = coordi%im_dim.x;
    skinY = coordi/im_dim.x;

        if( skinX >= 10 && skinY >= 10 )
    {
        skinX -= 10;
        skinY -= 10;
    }   

    unsigned int windowWidth = classifierMem[0];
    unsigned int windowHeight = classifierMem[1];   

    if(gx<srcWidth-windowWidth-1 && gy<srcHeight-windowHeight-1){
        bool stagePass = true;
        unsigned int index = 0;
        for(unsigned int i=numTotStage; i>0;i--){
            if(stagePass){
                unsigned int stageIndex = idxNumValStageArray[index++];
                float stageThres = classifierMem[stageIndex+2];
                float numFeatures = classifierMem[stageIndex+3];
                unsigned int featureIndex = stageIndex+4;
                float featureValue = 0.0;               

                for(unsigned int j=numFeatures; j>0;j--){
                    if(stagePass){
                        float featureThres=classifierMem[featureIndex++]*(windowWidth*windowHeight);
                        float succVal=classifierMem[featureIndex++];
                        float failVal=classifierMem[featureIndex++];
                        unsigned int numRegions = classifierMem[featureIndex++];
                        float regionValue =0.0;

                        for(unsigned int k=numRegions; k>0;k--){                    
                            float4 rectValue;
                            int4 regionP;

                            regionP.x=(int)(classifierMem[featureIndex])+skinX;
                            regionP.y=(int)(classifierMem[featureIndex+1])+skinY;
                            regionP.z=(int)(classifierMem[featureIndex+2])+regionP.x;
                            regionP.w=(int)(classifierMem[featureIndex+3])+regionP.y;

                            rectValue.x = read_imagef(input_image, sampler, regionP.xy).x;
                            rectValue.y = read_imagef(input_image, sampler, regionP.zy).x;
                            rectValue.z = read_imagef(input_image, sampler, regionP.xw).x;
                            rectValue.w = read_imagef(input_image, sampler, regionP.zw).x;

                            regionValue += dot(rectValue, (float4)(1.0f, -1.0f, -1.0f, 1.0f))*classifierMem[featureIndex+4];

                            featureIndex+=5;
                        }

                        featureValue += (regionValue < featureThres)?failVal:succVal;                   

                        if(j*2 == (unsigned int)numFeatures && featureValue*2 < stageThres) stagePass =false;
                    }
                }
                if(featureValue < stageThres)   stagePass =false;
                else if(index==numTotStage) write_imagef(output_image, (int2)(skinX, skinY), (0.1));
            }
        }   
    }else return;
}

время профилирования: исходная версия (без локальной памяти): 24 мс модифицированная версия (с локальной памятью): 28 мс

отредактировано: на самом деле localWorkSize NULL, потому что globalWorkSize всегда зависит от размера вектора, который помещает NDRangeKernel. Когда вы устанавливаете определенный localWorkSize, скорость обнаружения лиц падает ... Поэтому я попытался установить localWorkSize NUll, тогда скорость обнаружения лиц была хорошей. Так что я хочу, чтобы причина.

это код хоста:

    //localWorkSize[0] = 16;
    //localWorkSize[1] = 12; 
    numThreadsX=512;
    globalWorkSize[0] = numThreadsX;
    globalWorkSize[1] =  vecCoordinate.size()% numThreadsX == 0 ? vecCoordinate.size()/ numThreadsX :(vecCoordinate.size()/ numThreadsX) + 1;
    errNum = clEnqueueWriteBuffer(commandQueue,classifierMem,CL_TRUE,0,sizeof(float)*cntValArray,stageValArray,0,NULL,NULL); 
    errNum |= clEnqueueWriteBuffer(commandQueue,idxStageMem,CL_TRUE,0,sizeof(int)*haar.numStages,idxNumValStageArray,0,NULL,NULL); 
    errNum |= clSetKernelArg(kHaar_Cascade, 0, sizeof(memObjBuffer_Haar22), &memObjBuffer_Haar22);
    errNum |= clSetKernelArg(kHaar_Cascade, 1, sizeof(memObjBuffer22), &memObjBuffer22);
    errNum |= clSetKernelArg(kHaar_Cascade, 2, sizeof(cl_mem), &classifierMem);
    errNum |= clSetKernelArg(kHaar_Cascade, 3, sizeof(cl_mem), &idxStageMem);
    errNum |= clSetKernelArg(kHaar_Cascade, 4, sizeof(cl_int), &haar.numStages);
    errNum |= clSetKernelArg(kHaar_Cascade, 5, sizeof(cl_mem), &memVecCoordi);

    errNum = clEnqueueNDRangeKernel(commandQueue, kHaar_Cascade, 2, NULL,globalWorkSize, NULL,0, NULL, &event[3]);

android opencl viola-jones bank-conflict
person youngwan lee    schedule 01.05.2015    source источник
comment
На каком устройстве вы это запускаете?   -  person jprice    schedule 01.05.2015
comment
Затем уменьшите размер рабочей группы, чтобы увеличить максимальный локальный поток.   -  person huseyin tugrul buyukisik    schedule 01.05.2015
comment
@huseyintugrulbuyukisik Я не понимаю твоего комментария. Что значит создать максимальное локальное??   -  person youngwan lee    schedule 01.05.2015
comment
Если бы на каждую вычислительную единицу приходилось 256 потоков, при уменьшении ее до 64 объем доступной памяти на поток увеличивается в четыре раза.   -  person huseyin tugrul buyukisik    schedule 01.05.2015
comment
@huseyintugrulbuyukisik Моя проблема не в размере локальной памяти, но когда я использую больше локальной памяти, время профилирования медленнее, чем без использования локальной памяти, только глобальная версия памяти :( Когда я тестирую использование различных размеров локальной памяти, чем больше локальной памяти я использовал, тем медленнее время профилирования.   -  person youngwan lee    schedule 02.05.2015
comment
пожалуйста, укажите, насколько это медленнее - желательно также указать ms для каждой версии   -  person Jason Newton    schedule 03.05.2015
comment
@youngwanlee также указывает размеры рабочей группы   -  person Jason Newton    schedule 04.05.2015
comment
@youngwanlee, пожалуйста, не забудьте выбрать ответ и проголосовать за него. Кроме того, если вы не предоставите размеры рабочих групп, мы не сможем получить более конкретную информацию, но я уверен в приведенных мной причинах.   -  person Jason Newton    schedule 05.05.2015
comment
@JasonNewton на самом деле localWorkSize NULL, потому что globalWorkSize всегда зависит от размера вектора, который помещает NDRangeKernel. Когда вы устанавливаете определенный localWorkSize, скорость обнаружения лиц падает ... Поэтому я попытался установить localWorkSize NUll, тогда скорость обнаружения лиц была хорошей. Так что я хочу, чтобы причина.   -  person youngwan lee    schedule 05.05.2015
comment
вы можете определить локальный рабочий размер, распечатав из ядра с помощью printf   -  person Jason Newton    schedule 05.05.2015
comment
@JasonNewton Я не могу распечатать, потому что эта реализация на смартфоне   -  person youngwan lee    schedule 06.05.2015
comment
есть другие способы сделать это... сохранить его в другой глобальный буфер и распечатать после того, как ядро ​​установит его? в любом случае, хотя мне любопытно узнать, я не думаю, что ответ повлияет на вывод / ответы ниже.   -  person Jason Newton    schedule 06.05.2015
comment
@JasonNewton это ядро ​​​​Cascade рекомендует, что workGroupSize равен 192 (CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE) после печати localWorkSize, он показывает (128,1), (128,1) (8,17), (2,67), (2,67), (32, 5), (64,3),(64,3),(64,3),(128,1),(64,3),(64,3),(64,3),(128,1) ,(128,1),(128,1),(128,1),(128,1),(32,5),(128,1),(64,3),(128,1),( 128,1).....   -  person youngwan lee    schedule 07.05.2015

Ответы (2)


arrow_upward
0
arrow_downward

Есть множество причин:

  • Потоковое расхождение загрузки памяти
  • Барьеры не бесплатны
  • Дополнительные инструкции
  • На большинстве платформ постоянная память почти ничем не отличается от разделяемой памяти, поэтому дополнительная работа не приносит пользы.
  • Мы не знаем, сколько операций чтения экономит версия с общей памятью. Пожалуйста, сообщите нам размер рабочей группы. Чем он меньше, тем меньше операций чтения из нелокальной памяти мы сохраняем. Не то, чтобы это имело значение, если это постоянная память.

Профилируйте их обоих и посмотрите, показывают ли выходные данные профилировщика что-либо при максимальном использовании чего-либо, и укажите нам, что это такое.

Кроме того, я не уверен, что ваши localIdx и stage1Idx имеют смысл, они могут выходить за пределы массива и вызывать странное поведение. По крайней мере, для данного gx/gy вы выглядите так, как будто используете разные индексы из classifierMem.

person Jason Newton    schedule 03.05.2015

arrow_upward
0
arrow_downward

У вас есть пять значений с плавающей запятой (20 байт), которые считываются каждым рабочим элементом. Большинство современных графических процессоров будут иметь кэши, управляемые аппаратным обеспечением, которые находятся на том же уровне в иерархии памяти, что и локальная память, управляемая пользователем. В вашей версии, которая не использует локальную память, эти 20 байтов вполне благополучно будут храниться в кэше рядом с АЛУ и обеспечат все преимущества производительности повторного использования данных, на которые вы могли надеяться. Ваша версия, использующая локальную память, просто явно делает то же самое, но также добавляет кучу дополнительных накладных расходов на ручную инициацию копий и некоторую дополнительную синхронизацию между рабочими элементами, что неизбежно замедляет работу.

Локальная память может обеспечить повышение производительности в определенных сценариях при осторожном использовании. Однако во многих случаях аппаратные кэши работают лучше и оставляют вам гораздо более простой код.

person jprice    schedule 02.05.2015