我有以下开放文件:
kernel void ndft(
global float *re, global float *im, int num_values,
global float *spectrum_re, global float *spectrum_im,
global float *spectrum_abs,
global float *sin_array, global float *cos_array,
float sqrt_num_values_reciprocal)
{
// MATH MAGIC - DISREGARD FROM HERE -----------
float x;
float y;
float sum_re = 0;
float sum_im = 0;
size_t thread_id = get_global_id(0);
//size_t local_id = get_local_id(0);
// num_values = 24 (live environment), 48 (test)
for (int i = 0; i < num_values; i++)
{
x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal;
y = sin_array[thread_id * num_values + i] * sqrt_num_values_reciprocal;
sum_re = sum_re + re[i] * x + im[i] * y;
sum_im = sum_im - re[i] * y + x * im[i];
}
// MATH MAGIC DONE ----------------------------
//spectrum_re[thread_id] = sum_re;
//spectrum_im[thread_id] = sum_im;
//spectrum_abs[thread_id] = hypot(sum_re, sum_im);
float asdf = hypot(sum_re, sum_im); // this is just a dummy calculation
}与此类似,执行时间大约为15个(工作组规模=567个,14个工作组,共有7938个校对)。
然而,我当然需要收回行动的成果,这是最后几个方面的成果(结果)。 一旦我完成这些记忆作业(如果<代码>spectrum_X即为global,例如,或 local,则该盒子的退出时间将增至1.4至1.5毫升。
我认为,执行时间的增加是某种固定的间接费用,因此,我只是积累更多的数据,这样,由于这种效应而损失的相对时间将减少到最低限度。 但当我翻一番(即数据数额的两倍)时,执行时间也翻了一番(达到2.8 ~3.0 ms)。
我发现,即使我不满意这些线路的one,我的执行时间与我不满意的三者相同。 即使我添加了<代码>if (thread_id = 0)并加以操作,我也有同样的执行时间。 然而,这样做太慢了(我申请的上限是大约30个我们)。 当我以普通的《C法典》在我的《万国邮联》中加以管理时,它甚至预示着5倍的更快。
现在,我显然做了一些错误的事情,但我不敢肯定,在哪里开始寻求解决办法。
正如我对裁判所的答复所评论的那样,我也做了以下工作:
从上述法典来看,我把最后4条放在了上。
//spectrum_re[thread_id] = sum_re;
//spectrum_im[thread_id] = sum_im;
spectrum_abs[thread_id] = hypot(sum_re, sum_im);
//float asdf = hypot(sum_re, sum_im);As expected, execution time ~1.8 ms. The generated assembler code for my system is:
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Tue Apr 03 12:42:39 2012 (1333449759)
// Driver
//
.version 3.0
.target sm_21, texmode_independent
.address_size 32
.entry ndft(
.param .u32 .ptr .global .align 4 ndft_param_0,
.param .u32 .ptr .global .align 4 ndft_param_1,
.param .u32 ndft_param_2,
.param .u32 .ptr .global .align 4 ndft_param_3,
.param .u32 .ptr .global .align 4 ndft_param_4,
.param .u32 .ptr .global .align 4 ndft_param_5,
.param .u32 .ptr .global .align 4 ndft_param_6,
.param .u32 .ptr .global .align 4 ndft_param_7,
.param .f32 ndft_param_8
)
{
.reg .f32 %f;
.reg .pred %p;
.reg .s32 %r;
ld.param.u32 %r3, [ndft_param_2];
// inline asm
mov.u32 %r18, %envreg3;
// inline asm
// inline asm
mov.u32 %r19, %ntid.x;
// inline asm
// inline asm
mov.u32 %r20, %ctaid.x;
// inline asm
// inline asm
mov.u32 %r21, %tid.x;
// inline asm
add.s32 %r22, %r21, %r18;
mad.lo.s32 %r11, %r20, %r19, %r22;
setp.gt.s32 %p1, %r3, 0;
@%p1 bra BB0_2;
mov.f32 %f46, 0f00000000;
mov.f32 %f45, %f46;
bra.uni BB0_4;
BB0_2:
ld.param.u32 %r38, [ndft_param_2];
mul.lo.s32 %r27, %r38, %r11;
shl.b32 %r28, %r27, 2;
ld.param.u32 %r40, [ndft_param_6];
add.s32 %r12, %r40, %r28;
ld.param.u32 %r41, [ndft_param_7];
add.s32 %r13, %r41, %r28;
mov.f32 %f46, 0f00000000;
mov.f32 %f45, %f46;
mov.u32 %r43, 0;
mov.u32 %r42, %r43;
BB0_3:
add.s32 %r29, %r13, %r42;
ld.global.f32 %f18, [%r29];
ld.param.f32 %f44, [ndft_param_8];
mul.f32 %f19, %f18, %f44;
add.s32 %r30, %r12, %r42;
ld.global.f32 %f20, [%r30];
mul.f32 %f21, %f20, %f44;
ld.param.u32 %r35, [ndft_param_0];
add.s32 %r31, %r35, %r42;
ld.global.f32 %f22, [%r31];
fma.rn.f32 %f23, %f22, %f19, %f46;
ld.param.u32 %r36, [ndft_param_1];
add.s32 %r32, %r36, %r42;
ld.global.f32 %f24, [%r32];
fma.rn.f32 %f46, %f24, %f21, %f23;
neg.f32 %f25, %f22;
fma.rn.f32 %f26, %f25, %f21, %f45;
fma.rn.f32 %f45, %f24, %f19, %f26;
add.s32 %r42, %r42, 4;
add.s32 %r43, %r43, 1;
ld.param.u32 %r37, [ndft_param_2];
setp.lt.s32 %p2, %r43, %r37;
@%p2 bra BB0_3;
BB0_4:
// inline asm
abs.f32 %f27, %f46;
// inline asm
// inline asm
abs.f32 %f29, %f45;
// inline asm
setp.gt.f32 %p3, %f27, %f29;
selp.f32 %f8, %f29, %f27, %p3;
selp.f32 %f32, %f27, %f29, %p3;
// inline asm
abs.f32 %f31, %f32;
// inline asm
setp.gt.f32 %p4, %f31, 0f7E800000;
mov.f32 %f47, %f32;
@%p4 bra BB0_6;
mov.f32 %f48, %f8;
bra.uni BB0_7;
BB0_6:
mov.f32 %f33, 0f3E800000;
mul.rn.f32 %f10, %f8, %f33;
mul.rn.f32 %f47, %f32, %f33;
mov.f32 %f48, %f10;
BB0_7:
mov.f32 %f13, %f48;
// inline asm
div.approx.f32 %f34, %f13, %f47;
// inline asm
mul.rn.f32 %f39, %f34, %f34;
add.f32 %f38, %f39, 0f3F800000;
// inline asm
sqrt.approx.f32 %f37, %f38; // <-- this is part of hypot()
// inline asm
mul.rn.f32 %f40, %f32, %f37;
add.f32 %f41, %f32, %f8;
setp.eq.f32 %p5, %f32, 0f00000000;
selp.f32 %f42, %f41, %f40, %p5;
setp.eq.f32 %p6, %f32, 0f7F800000;
setp.eq.f32 %p7, %f8, 0f7F800000;
or.pred %p8, %p6, %p7;
selp.f32 %f43, 0f7F800000, %f42, %p8;
shl.b32 %r33, %r11, 2;
ld.param.u32 %r39, [ndft_param_5];
add.s32 %r34, %r39, %r33;
st.global.f32 [%r34], %f43; // <-- stores the hypot s result in spectrum_abs
ret;
}
确实,我的所有计算作业都是在座的,即:为<代码><<>hypot功能添加/组合的批次和<编码>。 根据上述法典,我删除了第二行:
st.global.f32 [%r34], %f43;
它是实际将数据储存在全球阵列spectrum_abs上的线。 然后,我使用<代码>clCreateProgramWithBinary,并将经过修改的编码文档用作投入。 处决时间已降至20个。