Linux 环境下 intel 核心显卡的 OpenCL 并行计算

zhujiang73

　 Intel CPU 内置的核心显卡也是可以做 OpenCL 并行计算的，虽然核心显卡的计算能力并不是很强．　

　 Linux 环境下有　Beignet　https://www.freedesktop.org/wiki/Software/Beignet/　

Beignet is an open source implementation of the OpenCL specification - a generic compute oriented API. This code base contains the code to run OpenCL programs on Intel GPUs which basically defines and implements the OpenCL host functions required to initialize the device, create the command queues, the kernels and the programs and run them on the GPU. The code base also contains the compiler part of the stack which is included in backend ...

    配合　Boost.Compute　https://github.com/boostorg/compute  和  OpenCV  可以处理图像，比如做一个边缘检测：

    参照  Boost.Compute 里这个例子 https://github.com/boostorg/comp ... cv_sobel_filter.cpp　效果大概是这样：



　下面的代码是在 GPU 平行的：

1. 
   
2. // Create sobel filter program
   
3. extern const std::string  str_cl_sobel = BOOST_COMPUTE_STRINGIZE_SOURCE
   
4. (
   
5.     //For out of boundary pixels, edge pixel
   
6.     // value is returned
   
7.     const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
   
8.    
   
9.     kernel void cl_kernel(read_only image2d_t src, write_only image2d_t dst)
   
10.     {
    
11.         int x = (int)get_global_id(0);
    
12.         int y = (int)get_global_id(1);
    
13. 
    
14.         if (x >= get_image_width(src) || y >= get_image_height(src))
    
15.                 return;
    
16. 
    
17.         //  [(x-1, y+1), (x, y+1), (x+1, y+1)]
    
18.         //  [(x-1, y  ), (x, y  ), (x+1, y  )]
    
19.         //  [(x-1, y-1), (x, y-1), (x+1, y-1)]
    
20. 
    
21.         //  [p02, p12,   p22]
    
22.         //  [p01, pixel, p21]
    
23.         //  [p00, p10,   p20]
    
24. 
    
25.         //Basically finding influence of neighbour pixels on current pixel
    
26.         float4 p00 = read_imagef(src, sampler, (int2)(x - 1, y - 1));
    
27.         float4 p10 = read_imagef(src, sampler, (int2)(x,     y - 1));
    
28.         float4 p20 = read_imagef(src, sampler, (int2)(x + 1, y - 1));
    
29. 
    
30.         float4 p01 = read_imagef(src, sampler, (int2)(x - 1, y));
    
31.         //pixel that we are working on
    
32.         float4 p21 = read_imagef(src, sampler, (int2)(x + 1, y));
    
33. 
    
34.         float4 p02 = read_imagef(src, sampler, (int2)(x - 1, y + 1));
    
35.         float4 p12 = read_imagef(src, sampler, (int2)(x,     y + 1));
    
36.         float4 p22 = read_imagef(src, sampler, (int2)(x + 1, y + 1));
    
37. 
    
38.         //Find Gx = kernel + 3x3 around current pixel
    
39.         //           Gx = [-1 0 +1]     [p02, p12,   p22]
    
40.         //                [-2 0 +2]  +  [p01, pixel, p21]
    
41.         //                [-1 0 +1]     [p00, p10,   p20]
    
42.         float3 gx = -p00.xyz + p20.xyz +
    
43.                     2.0f * (p21.xyz - p01.xyz)
    
44.                     -p02.xyz + p22.xyz;
    
45. 
    
46.         //Find Gy = kernel + 3x3 around current pixel
    
47.         //           Gy = [-1 -2 -1]     [p02, p12,   p22]
    
48.         //                [ 0  0  0]  +  [p01, pixel, p21]
    
49.         //                [+1 +2 +1]     [p00, p10,   p20]
    
50.         float3 gy = p00.xyz + p20.xyz +
    
51.                     2.0f * (- p12.xyz + p10.xyz) -
    
52.                     p02.xyz - p22.xyz;
    
53.         //Find G
    
54.         float3 g = native_sqrt(gx * gx + gy * gy);
    
55. 
    
56.         // we could also approximate this as g = fabs(gx) + fabs(gy)
    
57.         write_imagef(dst, (int2)(x, y), (float4)(g.x, g.y, g.z, 1.0f));
    
58.     }
    
59. );
    
60. 
    

复制代码

fender0107401

我也想玩这个。

zhujiang73

回复 2# fender0107401


    把能用的计算能力都用上．　

fender0107401

回复 3# zhujiang73

我经常要处理大量数据，一直都是靠多线程来实现并行，一直想弄这块。现在Boost里面有这个库了，应该能方便很多。

   

windoze

这玩意貌似有bug，在我的mac上老崩……

VIP_fuck

你们玩儿的东西真是高大上啊

流氓无产者

都高大尚啊，怎么那么多时间research啊

wlmqgzm

不错, 有时间的话, 也学习一下.