OpenCL 优化后的卷积代码

1217人阅读 评论(0) 收藏 举报
分类:

自己修改过的 可以直接运行 卷积模板 

void prefetch_texture_samples_8x4(image2d_t src, sampler_t s, __local float4 rgb [10][6], int2 gid, int2 lid)
{
  if (lid.x == 0) {
    // work-item 1 fetches all 60 rgb samples
    for (int i=-1; i<9; i++) {
      for (int j=-1; j<5; j++)
        rgb[i+1][j+1] = read_imagef(src, s, gid+(int2)(i, j));
    }
  }
  barrier(CLK_LOCAL_MEM_FENCE);
}

void prefetch_8x4_optimized(image2d_t src, sampler_t s, __local float4 rgb[10][6])
{
  // Coord of wi0 in NRDange
  int2 wi0Coord = (int2)(get_group_id(0)*8, get_group_id(1)*4); 

  // 2D to 1D address (from 8x4 to 32x1)
  int flatLocal = get_local_id(1)*8 + get_local_id(0); 

  // Only first 30 work-items load, each loads 2 values in sequence
  if (flatLocal < 30) 
  {
    /* Convert from flatLocal 1D id to 2D, 10x3 */
    int i = flatLocal % 10; // Width
    int j = flatLocal / 10; // Height
    
    /* 30 work iteams reads 10x3 values, 
     * values 0-9, 10-19, 20-29 from 10x6 - top half 
     */
    rgb[j][i] = read_imagef(src, s, (int2)(wi0Coord.x + i - 1, wi0Coord.y + j - 1));
    
    /* 30 work iteams reads 10x3 values, 
     * values 30-39, 40-49, 50-59 from 10x6 - bottom half 
     */
    rgb[j + 3][i] = read_imagef(src, s, (int2)(wi0Coord.x + i - 1, wi0Coord.y + j + 3 - 1));
  }
  barrier(CLK_LOCAL_MEM_FENCE);
}

__attribute__((reqd_work_group_size(8, 4, 1)))
__kernel void blur (image2d_t src, image2d_t dst, sampler_t s, float *weight)
{
  int2 gid = (int2)(get_group_id(0)*8, get_group_id(1)*4);
  int2 lid = (int2)(get_local_id(0),   get_local_id(1));
  float4 pixel = 0.0f;

  __local float4 rgb[10][6];
  prefetch_texture_samples_8x4(src, s, rgb, gid, lid);

  for (int j=-1; j<=1; j++) 
    for (int i=-1; i<=1; i++)
      pixel += rgb[lid.x+1+i][lid.y+1+i] * weight[(j+1)*3+i+1];
  
  int x = get_global_id(0);
  int y = get_global_id(1);  
  write_imagef(dst, (int2)(x, y), pixel/9.f);
}


查看评论

OpenCL编程:图像卷积

图像卷积,就是对图像所有像素进行一些特定的运算处理。这里涉及两个问题,一是读取图片文件信息,二是作何种卷积运算。第一个问题可在《freeimage存取图片数据》里找到答案。第二个问题可以baidu卷积...
  • jaccen
  • jaccen
  • 2016-05-13 17:35:19
  • 1772

关于《OPENCL异构并行计算》中卷积优化的分析

《OPENCL异构并行计算》中讲了如何利用OPENCL进行卷积运算,并给出了使用局部存储器优化的例子,这里对其进行简单分析...
  • qq_20028731
  • qq_20028731
  • 2017-04-16 17:13:39
  • 726

使用OpenCL+OpenCV实现图像卷积(一)

基于VS2010,使用OpenCL+OpenCV实现图像卷积处理。
  • icamera0
  • icamera0
  • 2017-06-08 07:50:00
  • 1175

OpenCL与CNN篇四:CNN从入门到使用

记录我从零到实现一个具体CNN网络中最有用的知识干货。 以细节为切入点,分享我对CNN网络的简洁。 本文致力于让你一篇文章理解CNN的具体实现与训练方法。 涉及理论不一一追述背景,主要讲解其如何应用。...
  • ShiAokai
  • ShiAokai
  • 2017-06-10 01:08:45
  • 1096

使用OpenCL+OpenCV实现图像卷积(二)

基于VS2010,使用OpenCL+OpenCV实现图像卷积处理。
  • icamera0
  • icamera0
  • 2017-06-08 07:59:34
  • 1551

OpenCL做并行滤波

本实验主要进行OpenCL一维信号的滤波;主要思路是以离散信号的序列点作为目标,一个工作项负责一个信号点的计算;这样做的好处是方便,相对于串行实现获得相当大的性能提升;但是每个工作项负载不均衡。...
  • u011028771
  • u011028771
  • 2016-10-15 11:25:31
  • 437

关于卷积

卷积卷积核大小(Kernel Size):定义了卷积操作的感受野。在二维卷积中,通常设置为3,即卷积核大小为3×3。步幅(Stride):定义了卷积核遍历图像时的步幅大小。其默认值通常设置为1,也可将...
  • xuke_2018
  • xuke_2018
  • 2018-04-15 22:43:01
  • 7

使用OpenCL+OpenCV实现图像卷积(三)

基于VS2010,使用OpenCL+OpenCV实现图像旋转功能。
  • icamera0
  • icamera0
  • 2017-06-08 08:05:45
  • 454

使用OpenCL+OpenCV实现图像旋转(二)

基于VS2010,使用OpenCL+OpenCV实现图像旋转功能。
  • icamera0
  • icamera0
  • 2017-05-10 22:37:37
  • 908

opencl:cl::make_kernel的进化

我之前的一篇博客《opencl:C++ 利用cl::make_kernel简化kernel执行代码》详细说明了如何使用OpenCL C++接口(cl.hpp)提供cl::make_kernel算子来简...
  • 10km
  • 10km
  • 2016-03-10 12:18:19
  • 1356
    个人资料
    持之以恒
    等级:
    访问量: 12万+
    积分: 2011
    排名: 2万+
    最新评论