将常量内存数组编译为 CUDA 中的立即值

Posted

技术标签:

【中文标题】将常量内存数组编译为 CUDA 中的立即值【英文标题】:compile constant memory array to immediate value in CUDA 【发布时间】:2015-01-20 09:16:55 【问题描述】:

我正在编写使用幂级数逼近函数的代码,并希望利用 #pragma unroll 和 FMA 指令,如下所示:

__constant__ double coeff[5] = 1.0,2.0,3.0,4.0,5.0; /* constant is fake here */

__device__ double some_function(double x) 
  double y;
  int i;
  y = coeff[0];
#pragma unroll
  for(i=1;i<5;i++) y = y*x + coeff[i];
  return y;

代码将被编译成这样的程序集:

ld.const.f64    %fd33, [coeff];
ld.const.f64    %fd34, [coeff+8];
fma.rn.f64      %fd35, %fd33, %fd32, %fd34;
ld.const.f64    %fd36, [coeff+16];
fma.rn.f64      %fd37, %fd35, %fd32, %fd36;
ld.const.f64    %fd38, [coeff+24];
fma.rn.f64      %fd39, %fd37, %fd32, %fd38;
ld.const.f64    %fd40, [coeff+32];
fma.rn.f64      %fd41, %fd39, %fd32, %fd40;

我想避免使用常量内存并像这样使用立即值:

mov.f64         %fd248, 0d3ED0EE258B7A8B04;
mov.f64         %fd249, 0d3EB1380B3AE80F1E;
fma.rn.f64      %fd250, %fd249, %fd247, %fd248;
mov.f64         %fd251, 0d3EF3B2669F02676F;
fma.rn.f64      %fd252, %fd250, %fd247, %fd251;
mov.f64         %fd253, 0d3F1745CBA9AB0956;
fma.rn.f64      %fd254, %fd252, %fd247, %fd253;
mov.f64         %fd255, 0d3F3C71C72D1B5154;
fma.rn.f64      %fd256, %fd254, %fd247, %fd255;
mov.f64         %fd257, 0d3F624924923BE72D;
fma.rn.f64      %fd258, %fd256, %fd247, %fd257;
mov.f64         %fd259, 0d3F8999999999A3C4;
fma.rn.f64      %fd260, %fd258, %fd247, %fd259;
mov.f64         %fd261, 0d3FB5555555555554;
fma.rn.f64      %fd262, %fd260, %fd247, %fd261;

我知道可以用#define宏来做,但是系数很多的时候很不方便。

是否有任何 C 数据类型修饰符(或编译器选项)可以将我的系数数组转换为立即值而不是使用常量内存?

我试过了,但它不适用于 static doublestatic __constant__ doublestatic __device__ double

我的最后一个问题是:我猜使用立即值应该比使用常量内存更快?

【问题讨论】:

对性能重要的是最终机器代码 (SASS),而不是中间表示 (PTX)。根据您的目标架构,FMA 指令可以直接引用常量内存,而无需单独的加载指令,这将达到最快速度(对常量内存的统一访问实际上与寄存器访问一样快)。您为 sm_35 编译的代码按以下顺序编译:DFMA R2, R4, c[0x3][0x0], R2; DFMA R2, R2, R4, c[0x3][0x10]; DFMA R2, R2, R4, c[0x3][0x18]; DFMA R4, R2, R4, c[0x3][0x20]; 即使性能一样,去掉它也会节省我恒定的内存空间(我的显卡只有64K) 【参考方案1】:

好的,您尝试做的事情是不可能的(至少不能使用 CUDA),这是因为 CUDA 禁止在全局范围内声明 static const 数组。 CUDA 要求将每个全局数组分配给特定的地址空间(__device____contant__ 等)。

但是有一些不便,这是可能的。

我收集了一些 SO 答案:

C++11: Compile Time Calculation of Array

Is it possible to develop static for loop in c++?

,请尊重那里的工作,并添加了一些CUDA。

你在这里:

您想要的是编译器为您完成繁琐的工作,因此您必须在编译时对所有内容进行评估:

首先我们需要一个静态数组来存储系数:

template <unsigned int index, long long... remPack> struct getValue;

template <unsigned int index, long long In, long long... remPack>
struct getValue<index, In, remPack...> 
  static const long long value = getValue<index - 1, remPack...>::value;
;

template <long long In, long long... remPack>
struct getValue<1, In, remPack...> 
  static const long long value = In;
;

template <long long... T> struct static_array 
  template <int idx> static __device__ int get()  return getValue<idx, T...>::value; 
;

static_array 在 C++ 类型系统中将值存储为 long long。稍后我会在答案中再次提到这一点。

接下来是必须展开的 for 循环。再次使用模板元编程:

template <int First, int Last, template <int> class Fn> struct static_for 
  __device__ double operator()(double x, double y) const 
      return static_for<First + 1, Last, Fn>()(x, Fn<First + 1>()(x, y));
  
;

template <int N, template <int> class Fn> struct static_for<N, N, Fn> 
  __device__ double operator()(double x, double y) const  return y; 
;

由于我们在编译时执行所有静态操作,因此有必要通过 operator() 的参数和返回表达式来移动每个“循环行程”的输入和输出。

这个解决方案是非常静态的,你可以通过更多的模板元编程来改进它。

好的,现在是有趣的部分。计算:

template <int i> struct Function 
  __device__ double operator()(double x, double y) 
    double c = __longlong_as_double(static_array<12, 34, 22, 55, 24>::get<i>());
    return y * x + c;
  
;

__device__ double some_function(double x) 
  return static_for<0, 5, Function>()(x, 0.0);

C++ 类型系统只允许整数类型作为非类型模板参数,因此我们必须将doubles 存储在long longs 中,然后使用CUDA 的__longlong_as_double() 函数将它们转换回来。这是我们目前必须接受的事情,并且可能会破坏您的交易,因为它不再“简单”了。但是,doublelong long 的转换器应该不难编写。

整个计算被包装在一个仿函数对象中,该对象从我们的static_loop 获取行程计数器作为模板参数。通过这个编译时间“行程计数器”,我们可以访问static_array 转换双回的long long 版本并计算FMA。

感谢 CUDA 编译器(这里做得很好)这是我使用 7.0 RC1 版本的 PTX 代码 (nvcc -ptx -arch=sm_35 test.cu):

.visible .func  (.param .b64 func_retval0) _Z13some_functiond(
        .param .b64 _Z13some_functiond_param_0
)

        .reg .f64       %fd<7>;

        ld.param.f64    %fd1, [_Z13some_functiond_param_0];
        fma.rn.f64      %fd2, %fd1, 0d0000000000000000, 0d000000000000000C;
        fma.rn.f64      %fd3, %fd2, %fd1, 0d0000000000000022;
        fma.rn.f64      %fd4, %fd3, %fd1, 0d0000000000000016;
        fma.rn.f64      %fd5, %fd4, %fd1, 0d0000000000000037;
        fma.rn.f64      %fd6, %fd5, %fd1, 0d0000000000000018;
        st.param.f64    [func_retval0+0], %fd6;
        ret;

【讨论】:

非常感谢您的努力。我认为它应该工作。但我对 C++ 元编程一无所知。最初我认为它应该通过添加一些编译器选项来工作。鉴于这种方法的复杂性。我会尝试写 20-30 行 __fma_rn(...) 来代替。【参考方案2】:

至少在 Cuda 8 中,本地 constexpr 数组可以正常工作,即对于展开的循环,*.ptx 包含立即值,而不是内存引用。示例(未经测试):

#define COEFF_VALUES  1.0, 2.0, 3.0, 4.0, 5.0 

__device__ double some_function( double x )

    constexpr double coeff[ 5 ] = COEFF_VALUES;
    double y;
    int i;
    y = coeff[ 0 ];
#pragma unroll
    for( i = 1; i < 5; i++ ) y = y*x + coeff[ i ];
    return y;

编译成这样的代码:

add.f64     %fd2, %fd1, 0d4000000000000000;
fma.rn.f64  %fd3, %fd1, %fd2, 0d4008000000000000;
fma.rn.f64  %fd4, %fd1, %fd3, 0d4010000000000000;
fma.rn.f64  %fd5, %fd1, %fd4, 0d4014000000000000;

【讨论】:

以上是关于将常量内存数组编译为 CUDA 中的立即值的主要内容,如果未能解决你的问题,请参考以下文章

CUDA学习5 常量内存与事件

CUDA 中的编译时信息

CUDA中常量内存的动态分配

如何用汇编语言读取内存值

CMake 是不是将包含的标头中的所有内容编译到可执行文件中,还是仅将主类中使用的部分编译为可执行文件?

Hbuilder中的less不能自动编译为css怎么办