CUDA 与 OpenCL 简明对比

又称『如何一把梭在 CUDA 和 OpenCL 之间进行转换』。

更多具体用法请参见最后的参考来源中的语言手册。

硬件层次

内存层次

CUDA 中，传一个指向 Global Memory 中地址的指针给 kernel, 只需像 C 语言一样写即可，如：kernel_func(float *a, ...). OpenCL 中进行同样的操作，需要在类型前面加上 __global, 即应该写成 kernel_func(__global float *a, ...).

CUDA 中，定义一个固定大小的 Shared Memory 数组可以在一个 kernel 函数体内用形如 __shared__ shm_a[1024]; 来定义。OpenCL 中定义一个固定大小的 Local Memory 可以在一个 kernel 函数体内用形如 __local local_a[1024]; 来定义。

所有在 CUDA kernel 内部定义和使用的变量都使用 Register 来保存。所有在 OpenCL kernel 内部使用的变量或者传入函数内部的参数都是 Private Memory。

线程层次

CUDA 中，每次启动 kernel 所有的 Thread 会排列成一个 Grid。每一个 Grid 中有若干个 Thread Block, 这些 Thread Block 排列成一个 3D 立方体：在 x, y, z 方向上，一个 Grid 有 gridDim.{x, y, z} 个 Thread Block, 每个 Thread Block 在 Grid 中 x, y, z 方向的编号为 blockIdx.{x, y, z}。每一个 Thread Block 中有若干个 Thread, 也排成一个 3D 立方体：在 x, y, z 方向上，一个 Thread Block 有 blockDim.{x, y, z} 个 Thread, 每个 Thread 在 Thread Block 中 x, y, z 方向的编号为 threadIdx.{x, y, z}.

OpenCL 中，每次启动 kernel 所有的 Work Item 会排列成一个工作空间（AMD 出的教程里面是这么叫的，没写英文）。每一个工作空间中有若干个 Work Group, 这些 Work Group 排列成一个 3D 立方体：在 x, y, z 方向上，一个工作空间有 G_{x, y, z} 个 Work Item，每个方向上有多少个 Work Group 则需要计算。每一个 Work Group 在工作空间中 x, y, z 方向的编号为 get_group_id({0, 1, 2}). 每一个 Work Group 中有若干个 Work Item, 也排成一个 3D 立方体：在 x, y, z 方向上，一个 Work Group 有 get_local_size({0, 1, 2}) 个 Work Item, 每个 Work Item 在 Work Group 中 x, y, z 方向的编号为 get_local_id({0, 1, 2}).

CUDA 中，每个 warp 中的 Thread 的执行步骤是一致的；在 OpenCL 中，每个 Wavefront 中的 Work Item 的执行步骤是一致的。或者说，应该将 CUDA 中的每个 warp 和 OpenCL 中的每个 Wavefront 视作一个 SIMT/SIMD 单元。

CUDA 中同步一个 Thread Block 里所有 Thread 需要使用 __syncthreads(). OpenCL 中使用 barrier(CLK_LOCAL_MEM_FENCE/CLK_GLOBAL_MEM_FENCE) 来同步同一个 Work Group 里所有 Work Item 在调用此函数之前对 Local/Global Memory 的读写操作。OpenCL 还可以使用 read_mem_fence(flags)/write_mem_fence(flag) 来同步在此之前对指定 flag （CLK_LOCAL_MEM_FENCE/CLK_GLOBAL_MEM_FENCE）的读/写操作。

定义函数

CUDA 中所有在设备上执行的函数需要在函数前面加上 __global__, __device__ 或 __host__ 关键字。使用 __global__ 关键字的函数即为从主机上启动在设备上执行的 kernel，其函数类型必须为 void. 使用 __device__ 关键字的函数只能在设备上从 kernel 中调用，不能在主机调用。如果不加关键字，即为默认使用 __host__, 函数只能在主机上调用。同一个文件中可以包含这三种函数，不需要分开写。

OpenCL 中定义设备上运行的 kernel 需要在函数前面加上 __kernel 关键字。如果需要从 kernel 中调用其它函数（而不是启动其他 kernel），被调用的函数前面不需要加上其他的关键字。kernel 函数必须写在同一个或几个文件中，并且这些文件中不应该包含在主机端执行的函数。

数据移动

CUDA 中分配设备上的内存和将数据在主机端和设备端移动的样例代码如下：

1
2
3
4
5
6
7
8
9
10
11

// 在设备端分配内存
cudaMalloc(void **devPtr, size_t size); 

// 在主机和设备之间拷贝数据, kind 指定操作类型
// cudaMemcpyHostToDevice:   主机端拷贝到设备端
// cudaMemcpyDeviceToDevice: 设备端拷贝到设备端
// cudaMemcpyDeviceToHost:   设备端拷贝到主机端
cudaMemcpy(void *dst, const void *src, size_t size, cudaMemcpyKind kind);

// 释放设备端内存
cudaFree(void *ptr);

同样功能的 OpenCL 代码就复杂很多了……需要写成下面这个样子：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37

cl_context       context;
cl_command_queue command_queue;
cl_int           err;

// 其他 OpenCL 初始化, context, command_queue 会被初始化

// 在设备端分配内存
// cl_mem buffer = clCreateBuffer(cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errcode);

// 将数据从主机端拷贝到设备端
err = clEnqueueWriteBuffer(
	cl_command_queue command_queue,
	cl_mem           buffer,
	cl_bool          blocking_write,
	size_t           buff_write_offset,
	size_t           buff_size_in_bytes,
	void             *host_src_ptr,
	cl_uint          num_events_in_wait_list,
	const cl_event   *event_wait_list,
	cl_event         *event
);

// 将设备从设备端拷贝到主机端
err = clEnqueueReadBuffer(
	cl_command_queue command_queue,
	cl_mem           buffer,
	cl_bool          blocking_read,
	size_t           buff_read_offset,
	size_t           buff_size_in_bytes,
	void             *host_dst_ptr,
	cl_uint          num_events_in_wait_list,
	const cl_event   *event_wait_list,
	cl_event         *event
)

// 释放设备端内存
clReleaseMemObject(cl_mem memObj);

使用其他设备内存如常量内存的方式请参见 API 手册。

启动方式

CUDA 中不使用 Stream 直接启动一个 kernel 的示例如下：

1
2
3
4
5

// 设置 kernel 运行的线程分组参数
dim3 grid(block_dim_x, block_dim_y, block_dim_z);
dim3 block(grid_dim_x, grid_dim_y, grid_dim_z);
// 启动 kernel
cuda_kernel<<<grid, block>>>(...);

OpenCL 中完成对应的操作实在太烦了。我按照网上的资料写了一个小的 Repo，在这里。

Reference

1. CUDA C Programming Guide
2. AMD OpenCL User Guide
3. OpenCL 1.2 Reference Pages
4. AMD上海研发中心编著的 OpenCL 教程
5. Experiences porting from CUDA to OpenCL
6. OpenCL 快速入门