将常量内存数组编译为 CUDA 中的立即值
compile constant memory array to immediate value in CUDA
我正在编写代码来使用幂级数来逼近函数,并且想利用#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 double
、static __constant__ double
和 static __device__ double
。
我的最后一个问题是:我想使用立即值应该比使用常量内存更快?
好的,按照您尝试的方式,您尝试做的事情是不可能的(至少对于 CUDA 是不可能的),这是因为 CUDA 禁止在全局范围内声明 static const
数组。 CUDA 要求将每个全局数组分配给特定地址 space(__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()
的 return 表达式移动每个 "loop trip" 的输入和输出。
这个解决方案是非常静态的,通过更多的模板元编程,您可以改进它。
好的,现在是有趣的部分。计算:
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 long
中,然后使用 CUDA 的 __longlong_as_double()
函数来转换它们背部。这是我们此时必须接受的事情,并且可能会破坏您的交易,因为它不再是 "simple"。然而,double
到 long long
的转换器不应该那么难写。
整个计算都包装在一个仿函数对象中,该对象从我们的 static_loop
中获取行程计数器作为模板参数。使用此编译时间 "trip counter" 我们可以访问 static_array
转换 long long
版本的双反并计算 FMA。
感谢 CUDA 编译器(它在这里做得非常好)这是 PTX 代码(nvcc -ptx -arch=sm_35 test.cu
)我使用的是 7.0 RC1 版本:
.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;
}
至少在 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;
我正在编写代码来使用幂级数来逼近函数,并且想利用#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 double
、static __constant__ double
和 static __device__ double
。
我的最后一个问题是:我想使用立即值应该比使用常量内存更快?
好的,按照您尝试的方式,您尝试做的事情是不可能的(至少对于 CUDA 是不可能的),这是因为 CUDA 禁止在全局范围内声明 static const
数组。 CUDA 要求将每个全局数组分配给特定地址 space(__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()
的 return 表达式移动每个 "loop trip" 的输入和输出。
这个解决方案是非常静态的,通过更多的模板元编程,您可以改进它。
好的,现在是有趣的部分。计算:
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 long
中,然后使用 CUDA 的 __longlong_as_double()
函数来转换它们背部。这是我们此时必须接受的事情,并且可能会破坏您的交易,因为它不再是 "simple"。然而,double
到 long long
的转换器不应该那么难写。
整个计算都包装在一个仿函数对象中,该对象从我们的 static_loop
中获取行程计数器作为模板参数。使用此编译时间 "trip counter" 我们可以访问 static_array
转换 long long
版本的双反并计算 FMA。
感谢 CUDA 编译器(它在这里做得非常好)这是 PTX 代码(nvcc -ptx -arch=sm_35 test.cu
)我使用的是 7.0 RC1 版本:
.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;
}
至少在 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;