OpenCL將數組從內存copy到顯存。本站提示廣大學習愛好者:(OpenCL將數組從內存copy到顯存)文章只能為提供參考,不一定能成為您想要的結果。以下是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