使用 OpenCL 加速 Web 应用程序

浏览器中的高性能计算

Web 计算语言 (WebCL) 支持 Web 应用程序在主机图像处理器上执行某些功能,极大地加速了诸多类型的数据处理例程,例如,数据分类、文本搜索和解决矩阵系统。通过本文您将了解 WebCL 的 JavaScript 类如何通过将计算任务交付到图像处理器来实现高速数字运算。

Matthew Scarpino, 软件开发人员, Eclipse Engineering, LLC

Matthew Scarpino 是位于加利福尼亚的 Mountain View 公司的一名软件开发人员,专门研究高性能图像处理。他是 OpenCL in Action 和 openclblog.com 的创始人。他目前正在编写一个可将 OpenGL 和 OpenCL 结合使用的开源固体建模工具。



2013 年 6 月 18 日

由于图像处理器 (GPU) 具有高度并行的基础架构,所以它可以比传统中央处理器 (CPU) 更快地执行某种类型的应用程序。开放计算语言 (OpenCL) 是最流行的语言之一,可驾驭 GPU 的强大功能。一个典型示例就是 Adobe® Premiere® Pro CS6,它通过在用户 GPU 上执行 OpenCL 例程来加速图像和视频的处理。

多家公司都认为 GPU 加速对基于浏览器的应用程序有很大的帮助,于是成立了一个工作组来促进这项技术的发展。在 2012 年 5 月,该工作组发布了 Web 计算语言 (WebCL) 的草案初稿。正如其主页(参阅 参考资料)所介绍的,WebCL 工作组旨在 “支持 Web 应用程序驾驭 Web 浏览器中的 GPU 和多核 CPU 处理”。

两家公司都已发布了相关扩展,支持从浏览器内部调用 WebCL 功能。Samsung 公司发布了 WebCL 来扩展 WebKit,一个可为 Mac OS X 上的 Apple Safari 浏览器提供强大功能的引擎。Nokia 发布了一个 Mozilla Firefox WebCL 插件,可在 Windows® 和 32 位 Linux® 操作系统上运行。由于具有广泛的开发人员基础,本文将重点介绍 Nokia 的实现。本文将会解释如何安装 WebCL,还会介绍 WebCL 编码的基础知识。然后示范如何使用 WebCL 在整个文本进行高速搜索。

安装 WebCL

WebCL 应用程序需要三个组件:OpenCL 软件开发工具包 (SDK)、Firefox 浏览器和 Nokia 的 Firefox 插件。

获得 OpenCL SDK

WebCL 应用程序在主机上调用 OpenCL 功能。因此,必须在运行 WebCL 应用程序之前安装 OpenCL。OpenCL 开发工具包是特定于设备和操作系统的;因此,要在带有 Nvidia GPU 的 Windows 计算机上运行例程,需要安装一个适用于 Windows 的 Nvidia OpenCL SDK。尽管本文不可能为所有 OpenCL SDK 安装都提供指南,但可为您指出正确的方向:

  • 在 AMD CPU 或 GPU 上执行例程,需要下载 AMD Accelerated Parallel Processing (APP) SDK(参阅 参考资料,获取相关链接)。
  • 在 Nvidia GPU 上执行例程,需要下载 Nvidia GPU Computing SDK(参阅 参考资料,获取相关链接)。
  • 在 Intel CPU 上执行例程,需要下载用于 OpenCL 的 Intel SDK(参阅 参考资料,获取相关链接)。

在一台计算机上可以安装多个 OpenCL SDK,并无冲突。

CPU 和 GPU

现代的 CPU 包括若干个被称为核心 的处理元素。每个核心都有自己的处理渠道和数据存储,并且核心之间的通信是使用类似直接访问内存的方法来实现的。为多核的 CPU 编程就好像领导一个训练有素的小型间谍团队:每个成员都能单独执行复杂的任务,但是整个流程作为团队运行良好。

相比之下,GPU 处理由工作项执行,这些工作项受限于内存和处理资源。工作项不擅长决策,对自身任务的协调更差。然而,GPU 的优势就在于它能同时执行成千上万的工作项。因此,为 GPU 编程就好像带领僵尸军团一样。尽管命令简单,但能完成大量工作。

在 Firefox 上安装 WebCL

安装完 OpenCL SDK 后,安装 WebCL 非常简单,步骤如下所示:

  1. 从 Mozilla 网站下载 Firefox(参阅 参考资料,获取相关链接)。
  2. 在 Firefox 中,登陆到 Nokia 的 WebCL 的网站(参阅 参考资料,获取相关链接)。单击 Nokia WebCL 扩展包
  3. 如有需要,可支持下载支持 Firefox 扩展包(一个 *.xpi 文件)。在 Software Installation 对话框中,单击 Install Now,然后重新启动浏览器。
  4. 为了测试安装,请登陆 Nokia 的 WebCL 网站,选择 Click here to check that you have WebCL enabled(单击此处查看您是否支持 WebCL)

    如果出现是一个对话框,Excellent! Your system does support WebCL(非常好!您的系统支持 WebCL),则表示 WebCL 就已成功安装。如果对话框显示,Unfortunately your system does not support WebCL(很抱歉,您的系统不支持 WebCL),那么您可能需要重新安装扩展包,或者在另一台计算机上访问 WebCL。


编写 WebCL 应用程序

WebCL 技术对新手来说会很神秘,所以牢记以下总目标非常重要:

  • 向设备递交一个函数。
  • 在设备上执行函数。
  • 从设备向主机传输输出

图 1 显示了这一过程。

图 1. WebCL 应用程序的操作
主机处理器 (CPU) 将一个内核部署到目标设备 (GPU)

WebCL 类和函数

要使用 Nokia 的工具集构建一个 WebCL 应用程序,则需要熟悉七个 JavaScript 类。表 1 列出每个类及其重要的函数。

表 1. 重要的 WebCL 类和函数
用途函数函数的描述
IWebCL包含访问平台和设备的静态函数getPlatformIDs()返回一个 IWebCLPlatform 对象数组
createContext(properties, devices)返回一个包含给定设备的 IWebCLContext 对象
createContextFromType(properties, device_type)返回一个包含给定类型的设备的 IWebCLContext 对象
IWebCLPlatform代表一个安装在主机上的 OpenCL SDKgetDeviceIDs(device_type)返回一个给定类型(可选的) IWebCLDevice 对象
IWebCLDevice代表一个符合 OpenCL 标准的物理设备getDeviceInfo(name)返回有关设备的信息
IWebCLContext管理程序、内存对象和命令队列createProgramWithSource(source_text)从源代码创建一个 IWebCLProgram 对象
createBuffer(flags, size)创建一个 IWebCLMemoryObject 对象来保存给定大小的数据
createCommandQueue(device, properties)创建一个 IWebCLCommandQueue 对象,将命令发送到给定设备
IWebCLProgram源代码包含一个或者多个内核函数buildProgram(devices, options)编译特定设备的程序源码
createKernel(kernel_name)从具有给定名称的函数创建一个 IWebCLKernel
IWebCLKernel特殊编码函数,能够通过一个 OpenCL 编译设备执行setKernelArg(index, value, type)从内存对象创建一个内核函数参数
IWebCLCommandQueue支持主机和设备间的通信enqueueReadBuffer(mem_object, blocking, offset, size, data_object, wait_list)将一个内存对象从设备读取到主机
enqueueWriteBuffer(mem_object, blocking, offset, size, data_object, wait_list)将一个内存对象从设备读取到主机
enqueueTask(kernel, wait_list)使用一个单一工作项,将一个 IWebCLKernel 执行对象插入队列
enqueueNDRangeKernel(kernel, dim, offset, global_size, local_size, wait_list)使用多个工作项,将一个 IWebCLKernel 执行对象插入队列

理解这些类如何相互配合使用需要花费一些时间,但是,一旦编码完成,只需极小改动就可以将一个 WebCL 应用程序复制粘贴到其他应用程序中。

一个简单示例

示例代码(可以从 下载 部分获取)中包含一个名为 simple.html 的文件,提供了完整的 WebCL 应用程序。JavaScript 代码和大多数 WebCL 应用程序代码一样,可以分为五个步骤:

  1. WebCL 应用程序首先创建了一个 IWebCLContext 对象来管理内核部署。WebCL.createContextFromType 函数带两个变量:一个包含大量数值的数组和一个用于识别环境中设备类型的值。以下代码展示了工作原理。
    清单 1. 创建了一个环境,并访问其设备
    var platforms = WebCL.getPlatformIDs();
    var ctx_props = [WebCL.CL_CONTEXT_PLATFORM, platforms[0]];
    var ctx = WebCL.createContextFromType(ctx_props, WebCL.CL_DEVICE_TYPE_GPU);
    var devices = ctx.getContextInfo(WebCL.CL_CONTEXT_DEVICES);

    CL_DEVICE_TYPE_GPU 参数表示这个环境只包含 GPU 设备。最后一行将环境设备放置在一个数组中,这个数组通常只包含一个元素:IWebCLDevice,代表主机的 GPU。

  2. 环境从源代码创建了一个 IWebCLProgram。应用程序然后对程序进行编译,从其中一个内核函数中创建一个 IWebCLKernel。这可以通过以下代码完成:
    清单 2. 编译了程序,并创建了一个内核
    var program_src = "__kernel void basic(__global float4* in_vec,     \
                                           __global float4* out_vec) {  \
                           out_vec[0] = in_vec[0];                      \
                       }";
    var program = ctx.createProgramWithSource(program_src);
    program.buildProgram([devices[0]], "");
    var kernel = program.createKernel("basic");

    了解程序内核 之间的差异非常重要。程序 是包含一个或者多个函数的一段代码。内核 代表程序代码中的一个单一函数。

  3. 内核可以由名为 basic 的函数创建,该函数接收两个参数:in_vecout_vec。这两个参数都存储在设备的寻址空间,所以应用将会创建两个 IWebCLMemoryObjectin_buffout_buff),并将其变为内核参数。以下代码显示了如何完成这一工作。
    清单 3. 创建了内存对象,并将其变为内核参数
    var in_buff = ctx.createBuffer(WebCL.CL_MEM_READ_ONLY, 16);
    var out_buff = ctx.createBuffer(WebCL.CL_MEM_WRITE_ONLY, 16);
    kernel.setKernelArg(0, in_buff);
    kernel.setKernelArg(1, out_buff);

    createBuffer 的参数设置了内存对象属性。第一个参数将会识别内核参数是提供输入(CL_MEM_READ_ONLY)还是接收输出(CL_MEM_WRITE_ONLY)。createBuffer 的第二个参数将会识别数据大小,以字节为单位。每个内存对象包含四个 4 字节的浮点数,因此其大小被设置为 16。

  4. 现在已经完成了内核配置,应用程序将创建一个 IWebCLCommandQueue,将指令从主机传输到内核。以下代码能够创建指令队列,将数据写入内核的输入参数(in_vec),然后通过调用 enqueueTask 加载内核。
    清单 4. 创建了指令队列,写入输入数据,并加载内核
    var queue = ctx.createCommandQueue(devices[0], 0);
    var in_data = new Float32Array([1.5, 2.5, 3.5, 4.5]);
    queue.enqueueWriteBuffer(in_buff, false, 0, 16, in_data, []);
    queue.enqueueTask(kernel, []);

    输入数据以 32 位浮点数组的形式提供。内核数据无须提供浮点类数,但它的值必须存储在相似类型的数组中,例如 Int32Array 或者 UInt16Array

  5. 为了验证内核是否被正确执行,应用程序将从内核第二个参数中读取数据,该参数作为 out_buff 内存对象提供。以下代码将数据从 out_buff 中读取到名为 out_data 的数组中,并在 Web 页面上显示数组的内容。
    清单 5. 从设备中读取输出数据,并将其显示在页面上
    out_data = new Float32Array(4);
    queue.enqueueReadBuffer(out_buff, true, 0, 16, out_data, []);      
    var output = document.getElementById("output");
    output.innerHTML = "Output: " + out_data[0] + ", " + out_data[1] + ", ";
    output.innerHTML += out_data[2] + ", " + out_data[3];

    这些代码通过访问 HTML 文档中称之为 output 的元素来显示输出数据。想要查看页面的 HTML 内容,请下载本文的示例代码,并在 Firefox 中打开 simple.html。

工作项和 enqueueNDRangeKernel

上述应用程序通过调用 enqueueTask 函数执行。在学习 WebCL 时这样没有问题,但是专业应用程序决不会使用这个函数,因为 enqueueTask 只使用了单一的线程执行内核,这有悖于大量使用 GPU 这类多并行设备的目的。

在 WebCL 中,处理线程被称为工作项。要理解如何使用它们,则需要参考 清单 6 中处理线程的嵌套循环。

清单 6. 处理线程
for(i=a; i<A; i++) {
   for(j=b; j<B; j++) {
      for(k=c; k<C; k++) {
         process_data(i, j, k);
      }
   }
}

process_data 函数有三个参数:ijk。每个参数都有不同的偏移量(abc),以及不同的最大范围(ABC)。因此,参数范围就是 A - aB - bC - c。OpenCL 术语将索引范围称之为大小,指标数量称之为维数。因此,该循环有三个维度,大小分别为 A - aB - bC - c

WebCL 的最大优势就是使用工作项执行函数迭代,例如 process_data,而不是时间密集型循环。工作项有不同维度(12 或者 3),它们每个都有惟一标识符,这些标识符都对应循环的特定迭代。如果工作项是使用二维执行的,那么一个工作项可能执行有惟一 ID (4, 5) 的内核,而另一个执行 ID 是 (5, 4) 的内核。请记住,这些工作项是并行执行的。

在 WebCL 中配置工作项,必须执行 enqueueNDRangeKernel,而不是 enqueueTask。这是 WebCL 编程接口中最重要的函数,它接受 6 个参数:

  • kernelIWebCLKernel 对应所要执行的函数。
  • dim:执行函数所使用的工作项维数。
  • offsets:包含工作项标识符初始值的数组。
  • global_sizes:包含各个维中工作项总数的数组。
  • local_sizes:包含工作组中各个维度的工作项数量的数组。
  • wait_list:包含 IWebCLEvent 结构的数组,由命令等待列表组成。

作为一个例子,可以考虑 清单 7 中的循环。

清单 7. 配置了一个 enqueueNDRangeKernel
for(i=5; i<50; i++) {
   for(j=6; j<60; j++) {
      for(k=7; k<70; k++) {
         process_data(i, j, k);
      }
   }
}

维数是 3,偏移量是 (5, 6, 7),工作项大小是 (45, 64, 73)。要使用 WebCL 执行 process_data,可以创建一个 IWebCLKernel 函数,并调用 enqueueNDRangeKernel 执行内核,如下所示:

queue.enqueueTask(kernel, 3, [5, 6, 7], [45, 64, 73], [], []);

最后两个参数被设置为空数组。虽然工作组和事件处理都是 WebCL 的优势,但它们不在本文的讨论范围之内。下一节将会介绍如何编写内核函数。


内核编码

WebCL 内核代表一个可在 GPU 或者其他适用的设备上执行的函数。总体结构如下:

__kernel void func_name(parameter_list) {
   ...lines of code...
}

内核声明以 __kernel 开头,其返回值是 void。这就意味着该函数的输入/输出数据必须由其参数列表提供。内核函数必须按照 C99 进行编码,但是开发人员可以访问许多 WebCL 特有的数据类型和函数。

内核的数据类型

内核可以访问 C 的基本数据类型,并使用向量数据类型更快速地处理数据。这些类型是以 typen 的形式提供的,type 是基本数据类型,n 是向量中该类型元素的数量(通常是 2、4、8 和 16)。

例如,一个 float4 就是一个包含 4 个浮点数的向量。如果设备支持 float4 操作,那么向量上的每个操作将会同时影响向量上的所有元素。如果不受支持,那么内核编译器将相关操作分散到受支持的设备上。

初始化一个向量和初始化一个数组类似,但使用的是圆括号(()),而不是波形括号({})。清单 8 中的代码显示了如何重新声明和初始化两个向量:int4char16

清单 8. 声明和初始化了两个向量
int4 vector_a = (int4)(1, 2, 3, 4);
char16 vector_b = (char16)('H', 'e', 'l', 'l', 'o',
    'P', 'r', 'o', 'g', 'r', 'a', 'm', 'm', 'e', 'r', '!');

运算符和数学函数

内核通常可以访问 C 中使用的所有运算符,这就意味着,如果 vec_avec_b 具有相同的向量类型,那么 vec_a * vec_b 会返回一个包含其乘积的向量。逻辑运算符也可用,而且 (5 > 3) 在内核中返回与正则函数相同的值。如果向量牵涉到关系运算,结果为 true 时,向量的元素会被设置为 -1,结果为 false 时,则设置为 0

如果 C 函数没有涉及到数学、逻辑或者比较运算,那么可能无法在内核中访问到它。但是内核几乎可以调用标准数学库中的所有函数,从 acossqrt。这些函数接受向量和标量。所以,要计算 vecfloat4 的绝对值,需要调用 fabs(vec),似乎 vec 只是另一个值。

工作项函数

当工作项执行内核时,它的第一个任务通常是确定自己的 ID。它可以访问内核的维数和执行内核的工作项总数。可以使用以下函数完成这项任务:

  • get_global_id(int i):在 i 维返回工作项 ID。
  • get_work_dim():返回执行内核的工作项的维数。。
  • get_global_size(int i):在 i 维中返回执行内核的工作项总数。

清单 9 中的内核函数显示了工作项函数在实践中的应用。

清单 9. 实践中的工作项函数
__kernel void example(__global float* in_array, __global float* out_array) {
   int id = get_global_id(0);
   int N = get_global_size(0);
   out_array[id] = in_array[id]/N; 
}

执行这一内核的工作项是一维的,在 in_arrayout_array 中,每个元素都有一个工作项。在执行内核时,每个工作项读取 in_array 的元素,用它除以工作项的总数,然后将结果存储到 out_array 中。简而言之,这就是 WebCL 内核的工作方式。下一节将会使用类似过程进行文本搜索。


使用 WebCL 搜索文本

许多 Web 应用程序需要尽快检测大文本集中的模式。这类应用程序对 GPU 来说非常完美,专门用于庞大数据的高强度处理。在 Amy Lowell 的叙事诗 “The Great Adventure of Max Breuck” 中,文本搜索项搜索四个词(thatverytimemore)。图 2 显示了其运行情况。

图 2. 使用 OpenCL 向量搜索文本
比较输入文本和模式

这一算法使用了 3 个文件来实现,这 3 个文件都包含在本文的可下载档案中:

  • text_search.html:包含 Web 页面的 HTML 和部署内核的 JavaScript。
  • text_search.cl:包含 Web 页面的 HTML 和部署内核的 JavaScript。
  • poem.txt:包含诗歌的原始文本(在开头和结尾都插入了额外的字符)。

每个工作项读取 16 个诗歌字符,并检查要搜索的四个词是否出现过。如果发现匹配项,工作项就会自动地增加四个计数器中的一个,将结果返回给主机。这个函数如 清单 10 所示。

清单 10. 增加文本搜索计数器的函数
		__kernel void text_search(__global char* text,
                          __global int* match_count) {

   char16 pattern = (char16)('t', 'h', 'a', 't', 'v', 'e', 'r', 'y', 
                             't', 'i', 'm', 'e', 'm', 'o', 'r', 'e');
   char16 test_vector, check_vector;

   /* load global text into private buffer */
   test_vector = vload16(0, text + get_global_id(0));

   /* compare test vector and pattern */
   check_vector = test_vector == pattern;

   /* Check for 'that,' 'very,' 'time,' and 'more' */
   if(all(check_vector.s0123))
      atomic_inc(match_count);
   if(all(check_vector.s4567))
      atomic_inc(match_count + 1);
   if(all(check_vector.s89AB))
      atomic_inc(match_count + 2);
   if(all(check_vector.sCDEF))
      atomic_inc(match_count + 3);
}

Firefox 的内核输出如 图 3 所示。

图 3. 文本搜索结果
Amy Lowell 诗歌的词语计数

使用 grep 功能对结果进行验证。例如,指令 grep -o time poem.txt 返回了词语 time 在文件 poem.txt 中出现的次数。

文本搜索主机应用程序生成 30,865 个工作项来执行内核。不是所有项都同时执行,具体情况取决于 GPU。这不是主要问题。只要一个工作项组完成执行,另一个组就会开始。因此,对 “我能执行多少个工作项?” 这一常见问题,答案是 ”您想要多少都可以“。


结束语:WebCL 的未来

当您考虑 WebCL 提供的功能有多强大时,可用的浏览器扩展寥寥无几,这看起来可能有些奇怪。问题在于安全性方面。内核执行可能会产生错误,比如无限循环或死锁。这对 CPU 来说不是什么大问题,但从我的经验来讲,GPU 错误会导致计算机死机。使用 Ctrl-C 和 Ctrl-Alt-Delete 无法解决问题,必须重启系统。

为了解决这一问题,GPU 制造商保证在其产品中添加了更好的安全性。这意味着用户可以中断或者中止 GPU 的执行,而不会导致系统崩溃。如果这是常见用法,WebCL 将会从古怪的东西变为必需的东西。通过 WebCL 的高速执行,浏览器应用能够提供与台式计算机应用程序相同的性能来处理数据和显示图像。


下载

描述名字大小
WebCL 开发的样例代码webcl_src.zip10KB

参考资料

学习

获得产品和技术

讨论

  • developerWorks 社区:探索由开发人员推动的博客、论坛、组和维基,并与其他 developerWorks 用户进行交流。

条评论

developerWorks: 登录

标有星(*)号的字段是必填字段。


需要一个 IBM ID?
忘记 IBM ID?


忘记密码?
更改您的密码

单击提交则表示您同意developerWorks 的条款和条件。 查看条款和条件

 


在您首次登录 developerWorks 时,会为您创建一份个人概要。您的个人概要中的信息(您的姓名、国家/地区,以及公司名称)是公开显示的,而且会随着您发布的任何内容一起显示,除非您选择隐藏您的公司名称。您可以随时更新您的 IBM 帐户。

所有提交的信息确保安全。

选择您的昵称



当您初次登录到 developerWorks 时,将会为您创建一份概要信息,您需要指定一个昵称。您的昵称将和您在 developerWorks 发布的内容显示在一起。

昵称长度在 3 至 31 个字符之间。 您的昵称在 developerWorks 社区中必须是唯一的,并且出于隐私保护的原因,不能是您的电子邮件地址。

标有星(*)号的字段是必填字段。

(昵称长度在 3 至 31 个字符之间)

单击提交则表示您同意developerWorks 的条款和条件。 查看条款和条件.

 


所有提交的信息确保安全。


static.content.url=http://www.ibm.com/developerworks/js/artrating/
SITE_ID=10
Zone=Web development, Open source
ArticleID=930781
ArticleTitle=使用 OpenCL 加速 Web 应用程序
publish-date=06182013