本文介绍一种新的CUDA编程范式:NVRTC+动态实例化。这种范式相对于传统的基于CUDA Runtime的编程的强大之处在于,在提供和模板编程同样的灵活性的同时,得到的库可以在各种不同的编程语言中使用,只要该语言可以和C互操作。
本文的英文版本:Programming GPU across the Language Boundariesfynv.github.io
ThrustRTC 这个GPU计算库就是基于这个范式开发的。
模板在GPU编程中的重要性和局限性
计算密集型的库基本上都是用C/C++或者其他编译型的语言写的。通常情况下,在其他语言中,包括解释型语言中也可以使用这些库。有一部分现存的GPU计算库也是这样的,只要这些库的所有接口都是非模板的Host端API,这些GPU库和普通的CPU库在使用上并没有什么区别。
但是,也有一些GPU库不是这样的。这些GPU库使用了模板技术,以源代码形式发布,在被使用之前不编译。著名的 Thrust 就是这样的。模板技术在GPU编程中尤其不可或缺,因为(动态)多态技术在GPU上的效率比较低下,为了多态特性,常常需要用模板编程来弥补。对于库来说,多态性常常显得尤其重要,因为我们总是试图让一个库具有尽可能广的适用范围。
然而模板库这种“未编译”的形式导致它们只能在特定的语言中使用。
template
void thrust::replace(ForwardIterator first,
ForwardIterator last,
const T & old_value,
const T & new_value)
这个是 Thrust 里一个简单的函数。所有的参数都模板化了&