Почему запись в память занимает больше времени, чем чтение в графическом процессоре? - PullRequest
Новая
       103

Почему запись в память занимает больше времени, чем чтение в графическом процессоре?

0 голосов
/ 19 мая 2018

Вот часть моего кода, где я обновляю некоторые значения и записываю его в другую часть моей памяти 

__kernel void update(
        __global float4 *f,
        __global float4 *p,
        const unsigned int n,
        __local float4 *sharedP){

        float4 val = (float4)(0.f,0.f,0.f,0.f);

        size_t lID = get_local_id(0);
        size_t gID = get_global_id(0);

        if(gID < n){
            float4 p1 = p[gID];
            for(size_t i = 0; i < n; ++i){
                if(i!=gID){
                    float4 p2 = p1 - p[i];
                    //Some other calculations to finally update 'val' 
                }
            }
            f[gID] = val;
         }
}

Здесь каждый поток читает из глобальной памяти n + 1 раз, но записывает обратно только один раз(Первоначально я думал, что узким местом было повторное обращение к глобальной памяти, поэтому я внес некоторые изменения, чтобы скопировать блок данных в общую память, поменять его местами и другие вещи. Даже это не улучшило производительностьмного).Однако позже я обнаружил, что узкое место связано с f[gID] = val;.Например, когда ядро ​​запускается 200 раз с n = 8192, это занимает 3.5s;но если я закомментирую f[gID] = val;, это займет всего 0.05s.

Код для другого метода 

__kernel void update(
        __global float4 *f,
        __global float4 *p,
        const unsigned int n,
        __local float4 *sharedP){

        float4 val = (float4)(0.f,0.f,0.f,0.f);

        size_t lID = get_local_id(0);
        size_t gID = get_global_id(0);

        if(gID < n){
            sharedP[lID] = p[gID];
            barrier(CLK_LOCAL_MEM_FENCE);
            float4 p1 = sharedP[lID];

            //To prevent i!=gID of other method
            for(size_t i = 0; i < get_local_size(0);; ++i){
                if(lID!=i){
                    float4 p2 = p1 - sharedP[i];
                    //Some other calculations to update 'val' at 0
                }
            }

            for(size_t block = 1; block < (size_t)(get_global_size(0)/get_local_size(0)); ++block){
                size_t idx = lID + (get_group_id(0) + block ) * get_local_size(0);

                if( idx >= n)
                    idx -= n;

                sharedP[lID] = p[idx];
                barrier(CLK_LOCAL_MEM_FENCE);
                for(size_t i = 0; i < get_local_size(0); ++i){
                    float4 p = p1 - sharedP[i];
                    //Some other calculations to update 'val'
                }   
                f[gID] = val;
            }
        }
}

Так как же запись в память заставляет ее так сильно замедляться и почемуиспользование разделяемой памяти, не обеспечивающее значительного увеличения производительности?

...