OpenCL將數組從內存copy到顯存

時間 2019-12-09

標籤 opencl 數組內存 copy 简体版

原文原文鏈接

原本想對上一篇博客作優化，優化效果不明顯。但知識點仍是要記一下。函數

初衷是想把上一篇博客中定義域的計算搬到CPU來計算，由於定義域的計算對於每個kernel都是同樣的，因此直接讀取應該是能夠進一步減少kernel的執行時間的。優化

個人思路的初衷是將這塊的數據送到顯存以後再送到寄存器中，從寄存器讀取的時間應該是很快的，經過這樣把計算的時間改成讀取的時間。固然，讀取寄存器的時間是否比計算更短，這個確實應該質疑，可是對於比較複雜的計算，我以爲直接讀應該是比計算更快的。而對於這部分數據，CPU計算應該會比GPU更快。固然，還應當考慮數據量的大小，從內存搬到顯存也是須要時間的。spa

1.C++代碼

..................

int ksize = 11;
float sigma_d = 3.0;
float *dkl = new float[ksize*ksize];
for (int i = -ksize/2; i <= ksize/2; i++){
    for (int j = -ksize/2; j <= ksize/2; j++){
        dkl[(i+ksize/2)*ksize + (j+ksize/2)] = -(i*i + j*j) / (2 * sigma_d*sigma_d);
    }
}

cl_mem d_dkl;
d_dkl = clCreateBuffer(context, CL_MEM_READ_ONLY, ksize*ksize*sizeof(float), NULL,NULL);
clEnqueueWriteBuffer(commandQueue, d_dkl, CL_TRUE, 0, ksize*ksize*sizeof(float), dkl, 0, NULL, NULL);

........................


errNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_dkl);
errNum |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &ksize);

........................

delete[] dkl;

...................

主要就是clCreateBuffer函數和clEnqueueWriteBuffer函數的用法。.net

2.kernel代碼

const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

kernel void bilateralBlur(read_only image2d_t src, write_only image2d_t dst, __constant float* dkl, int ksize)  
{
    int x = (int)get_global_id(0);  
    int y = (int)get_global_id(1);  
    if (x >= get_image_width(src) || y >= get_image_height(src))  
        return;  

    float sigma_d = 3.0;
    float sigma_r = 0.1;

    float4 fij = read_imagef(src, sampler, (int2)(x, y));

    float alpha = 0.2;
    float4 fkl;
    float4 rkl;
    float4 wkl;

    int index = 0;

    float4 numerator = (float4)(0.0f,0.0f,0.0f,0.0f);
    float4 denominator = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
    for (int K = -ksize / 2; K <= ksize / 2; K++)
    {
        for (int L = -ksize / 2; L <= ksize / 2; L++)
        {
            fkl = read_imagef(src, sampler, (int2)(x + K, y + L));

            rkl.x = -(fij.x - fkl.x)*(fij.x - fkl.x) / (2 * sigma_r*sigma_r);
            rkl.y = -(fij.y - fkl.y)*(fij.y - fkl.y) / (2 * sigma_r*sigma_r);
            rkl.z = -(fij.z - fkl.z)*(fij.z - fkl.z) / (2 * sigma_r*sigma_r);

            wkl.x = exp(-dkl[index] + rkl.x);
            wkl.y = exp(-dkl[index] + rkl.y);
            wkl.z = exp(-dkl[index] + rkl.z);
            index++;

            numerator.x += fkl.x * wkl.x;
            numerator.y += fkl.y * wkl.y;
            numerator.z += fkl.z * wkl.z;

            denominator.x += wkl.x;
            denominator.y += wkl.y;
            denominator.z += wkl.z;
        }
    }
    
    float4 gij = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
    if (denominator.x > 0 && denominator.y > 0 && denominator.z)
    {
        gij.x = numerator.x / denominator.x;
        gij.y = numerator.y / denominator.y;
        gij.z = numerator.z / denominator.z;

        gij.x = fij.x*alpha + gij.x*(1.0 - alpha);
        gij.y = fij.y*alpha + gij.y*(1.0 - alpha);
        gij.z = fij.z*alpha + gij.z*(1.0 - alpha);
    }

    write_imagef(dst, (int2)(x, y), gij);
}