OpenCL 2.0 异构计算 [第三版] (中文)
  • Introduction
  • 序言
  • 第1章 简介异构计算
    • 1.1 关于异构计算
    • 1.2 本书目的
    • 1.3 并行思想
    • 1.4 并发和并行编程模型
    • 1.5 线程和共享内存
    • 1.6 消息通讯机制
    • 1.7 并行性的粒度
    • 1.8 使用OpenCL进行异构计算
    • 1.9 本书结构
  • 第2章 设备架构
    • 2.1 介绍
    • 2.2 硬件的权衡
    • 2.3 架构设计空间
    • 2.4 本章总结
  • 第3章 介绍OpenCL
    • 3.1 简介OpenCL
    • 3.2 OpenCL平台模型
    • 3.3 OpenCL执行模型
    • 3.4 内核和OpenCL编程模型
    • 3.5 OpenCL内存模型
    • 3.6 OpenCL运行时(例子)
    • 3.7 OpenCL C++ Wapper向量加法
    • 3.8 CUDA编程者使用OpenCL的注意事项
  • 第4章 OpenCL案例
    • 4.1 OpenCL实例
    • 4.2 直方图
    • 4.3 图像旋转
    • 4.4 图像卷积
    • 4.5 生产者-消费者
    • 4.6 基本功能函数
    • 4.7 本章总结
  • 第5章 OpenCL运行时和并发模型
    • 5.1 命令和排队模型
    • 5.2 多命令队列
    • 5.3 内核执行域:工作项、工作组和NDRange
    • 5.4 原生和内置内核
    • 5.5 设备端排队
    • 5.6 本章总结
  • 第6章 OpenCL主机端内存模型
    • 6.1 内存对象
    • 6.2 内存管理
    • 6.3 共享虚拟内存
    • 6.4 本章总结
  • 第7章 OpenCL设备端内存模型
    • 7.1 同步和交互
    • 7.2 全局内存
    • 7.3 常量内存
    • 7.4 局部内存
    • 7.5 私有内存
    • 7.6 统一地址空间
    • 7.7 内存序
    • 7.8 本章总结
  • 第8章 异构系统下解析OpenCL
    • 8.1 AMD FX-8350 CPU
    • 8.2 AMD RADEON R9 290X CPU
    • 8.3 OpenCL内存性能的考量
    • 8.4 本章总结
  • 第9章 案例分析:图像聚类
    • 9.1 图像聚类简介
    • 9.2 直方图的特性——CPU实现
    • 9.3 OpenCL实现
    • 9.4 性能分析
    • 9.5 本章总结
  • 第10章 OpenCL的分析和调试
    • 10.1 设置本章的原因
    • 10.2 使用事件分析OpenCL代码
    • 10.3 AMD CodeXL
    • 10.4 如何使用AMD CodeXL
    • 10.5 使用CodeXL分析内核
    • 10.6 使用CodeXL调试OpenCL内核
    • 10.7 使用printf调试
    • 10.8 本章总结
  • 第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度
    • 11.1 简要介绍现状
    • 11.2 简单介绍C++ AMP
    • 11.3 编译器的目标 —— OpenCL 2.0
    • 11.4 C++ AMP与OpenCL对比
    • 11.5 C++ AMP的编译流
    • 11.6 编译之后的C++ AMP代码
    • 11.7 OpenCL 2.0提出共享虚拟内存的原因
    • 11.8 编译器怎样支持C++ AMP的线程块划分
    • 11.9 地址空间的推断
    • 11.10 优化数据搬运
    • 11.11 完整例子:二项式
    • 11.12 初步结果
    • 11.13 本章总结
  • 第12章 WebCL:使用OpenCL加速Web应用
    • 12.1 介绍WebCL
    • 12.2 如何使用WebCL编程
    • 12.3 同步机制
    • 12.4 WebCL的交互性
    • 12.5 应用实例
    • 12.6 增强安全性
    • 12.7 服务器端使用WebCL
    • 12.8 WebCL的状态和特性
  • 第13章 其他高级语言中OpenCL的使用
    • 13.1 本章简介
    • 13.2 越过C和C++
    • 13.3 Haskell中使用OpenCL
    • 13.4 本章总结
Powered by GitBook
On this page

Was this helpful?

  1. 第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度

11.9 地址空间的推断

上节中声明的每个OpenCL变量都具有自己地址空间限定符,用来区分这个变量是在哪端内存区域上分配的。地址空间对于OpenCL来说是十分中要的特性。将数据放入不同的内存区域,OpenCL程序在获得高性能的同时,保证了数据一致性。这个特性通常不会出现一些比较高级的语言中,比如C++ AMP。高级语言将数据放入通用地址空间内,从而就不用显式的说明这些内存是在哪里开辟的。OpenCL中声明的变量如果没有限定符,那么默认在私有内存上进行分配,这就违反了C++ AMP中的既定行为。举个例子,如果将tile_static的声明限定于私有,那么这个对象上的数据将不会与其他工作项共享,并且计算得到的结果是错误的。为了解决这个矛盾,就需要为每个声明和内存访问添加正确的地址空间信息。

CLamp中,生成OpenCL位码之后,需要在通过一次LLVM的转换,为相应的变量声明添加上正确的地址空间信息。理论上每个声明进行地址空间的推断是不可行的,因为分析器看不到整个程序,所以无法判断哪些内核要和哪些变量进行交互。不过,实际使用的程序中,推断地址空间是可行的。

array和array_view的实现都为推断地址空间提供着线索。C++ AMP中,只有通过array和array_view才能将大量的数据传入内核。C++ AMP运行时为内核的参数列表预留了指针。内核在使用这些数据时,只需要访问相关的指针即可。这些指针都会描述成全局的,因为这些数据时要每个工作项都可见的。推断过程的依据就是内核函数的参数列表,相关指针限定为全局,并且通过这些指针对更新所有内存操作。

tile_static数据的声明不能通过模式分析进行判别,所以CLamp的前端编译器要保存这些声明。当前的CLamp实现中,限定符声明tile_static的部分,使用一段特殊的位码进行表示。推断过程会将tile_static属性传递给任意一个指针,这些指针能获取这些变量的地址,然后将其添加到对应的OpenCL声明中。

我们看一个简单的C++ AMP实例,通过这个实例我们来了解转换是如何进行的:

void mm_kernel(int *p, int n){
  tile_static int tmp[30];
  int id = get_global_id(0);
  tmp[id] = 5566;
  barrier(0);
  p[id] = tmp[id];
}

通过CLamp初始化之后,代码将完全转化成LLVM IR。这个阶段中,地址空间是缺失的,并且这段代码会产生一个不正确的结果。注意变量tmp会放在一个特殊的ELF字段中(“clamp_opencl_local”):

@mm_kernel.tmp = internal unnamed_addr global[30xi32] zeroinitializer, align 16, section "clamp_opencl_local"

define void @mm_kernel(i32 *nocapture %p, i32 %n){
  %1 = tail call i32 bitcast (i32(...)* @get_global_id to i32 (i32)*)(i32 0)
  %2 = sext i32 % to i64
  %3 = getelementptr inbounds[30 x i32]* @mm_kernel.tmp, i64 0, i64 %2
  %4 = tail call i32 bitcast (i32(...)* @barrier to i32(i32)*)(i32 0) #2
  %5 = load i32 *%3, align 4, !tbaa!1
  %6 = getelementptr inbounds i32* %p, i64 %2
  store i32 %5, i32 * %6, align 4, !tbaa !1
  ret void
}

CLamp分析完成后,正确的地址空间信息就添加到对应的声明中去(mm_kernel.tmp中的一些内存操作)。正确的LLVM IR如下所示:

@mm_kernel.tmp = internal addrspace(3) unnamed_addr global[30xi32] zeroinitializer, align 4

define void @mm_kernel(i32 addrspace(1)*nocapture %p, i32 %n){
  %1 = tail call i32 bitcast(i32 (...)* @get_global_id to i32(i32)*)(i32 0)
  %2 = getelementptr inbounds[30 x i32] addrspace(3)* @mm_kenrel.tmp, i32 0, i32 %1 store i32 5566, i32 addrspace(3) %2, align4, !tbaa!2
  %3 = tail call i32 bitcast (i32(...) * @barrier to i32(i32)*)(i32 0)
  %4 = load i32 addrspace(3)* %2, align4, !tbaa !2
  %5 = getelementptr inbounds i32 addrspace(1) * %p, i32 %1
  store i32 %4, i32 addrspace(1)* %5, align 4, !tbaa !2
  ret void
}
Previous11.8 编译器怎样支持C++ AMP的线程块划分Next11.10 优化数据搬运

Last updated 6 years ago

Was this helpful?