私は次のOpenCLカーネルを持っています：
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 usです（ワークグループのサイズ= 567、14のワークグループ、合計7938スレッド）。
ただし、もちろん、最後の数行の目的である操作の結果を取得する必要があります（コメントアウト）。これらのメモリ操作の1つを実行するとすぐに（例のように、またはであるかどうかspectrum_Xは関係ありません）、カーネルの実行時間は約1.4〜1.5ミリ秒に増加します。globallocal
実行時間の増加はある種の固定オーバーヘッドだと思ったので、より多くのデータを蓄積するだけで、その影響によって失われる相対的な時間が最小限に抑えられます。しかし、スレッド数を2倍にすると（つまり、データ量が2倍になると）、実行時間も2倍になります（2.8〜3.0ミリ秒）。
これらの行の1つだけをコメント解除しても、3つすべてをコメント解除した場合と同じ実行時間が得られることがわかりました。を追加しif (thread_id == 0)て実行しても、実行時間は同じです。ただし、この方法では遅すぎます（私のアプリケーションの上限は約30 usです）。CPU上の通常の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);予想どおり、実行時間は約1.8ミリ秒です。私のシステム用に生成されたアセンブラコードは次のとおりです。
//
// 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;
}
sqrt確かに、私の計算操作はすべてそこにあります-関数のためだけでなく、たくさんの追加/マルチがありhypotます。上記のasmコードから、最後から2番目の行を削除しました。
st.global.f32 [%r34], %f43;
これは、実際にデータをグローバル配列に格納する行ですspectrum_abs。次にclCreateProgramWithBinary、変更したasmコードファイルを入力として使用しました。実行時間は20usに減少しました。