OpenCL将数组从内存copy到显存

本来想对上一篇博客做优化，优化效果不明显。但知识点还是要记一下。

初衷是想把上一篇博客中定义域的计算搬到CPU来计算，因为定义域的计算对于每一个kernel都是一样的，所以直接读取应该是可以进一步减小kernel的执行时间的。

我的思路的初衷是将这块的数据送到显存之后再送到寄存器中，从寄存器读取的时间应该是很快的，通过这样把计算的时间改为读取的时间。当然，读取寄存器的时间是否比计算更短，这个确实应该质疑，但是对于比较复杂的计算，我觉得直接读应该是比计算更快的。而对于这部分数据，CPU计算应该会比GPU更快。当然，还应当考虑数据量的大小，从内存搬到显存也是需要时间的。

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函数的用法。

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);
}

与上一博客的代码相比，主要就是把dkl的计算改为了读取，ksize也通过参数传进来。

3.结果

与上一篇3.42ms相比，有零点几毫秒的优化。不过考虑CPU的计算，优化应该更小，或者没有，或者稍差。

当然，我这里的计算简单，对于复杂的计算，应该还是可以考虑这种优化方法的。

下一步考虑内存优化，增大粒度。

代码：http://download.csdn.net/download/qq_33892166/9771206

转载于:https://www.cnblogs.com/betterwgo/p/6506844.html