OpenCL でより効率的に実行するために、2D 配列の多くの重複しているがオフセットされたブロックに対する操作をどのように構造化できますか?
たとえば、次の OpenCL カーネルがあります。
__kernel void test_kernel(
read_only image2d_t src,
write_only image2d_t dest,
const int width,
const int height
)
{
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
int2 pos = (int2)(get_global_id(0), get_global_id(1));
int2 pos0 = (int2)(pos.x - pos.x % 16, pos.y - pos.y % 16);
uint4 diff = (uint4)(0, 0, 0, 0);
for (int i=0; i<16; i++)
{
for (int j=0; j<16; j++)
{
diff += read_imageui(src, sampler, (int2)(pos0.x + i, pos0.y + j)) -
read_imageui(src, sampler, (int2)(pos.x + i, pos.y + j));
}
}
write_imageui(dest, pos, diff);
}
正しい結果が得られますが、遅いです... 1k x 1k 入力の NVS4200M でわずか 25 GFLOPS です。(ハードウェア仕様は 155 GFLOPS です)。これは、メモリアクセスパターンに関係していると思います。各作業項目は、16x16 領域内のすべての隣接ブロックと同じ 16x16 データ ブロックを 1 つ読み取ります。また、データの別のオフセット ブロックは、ほとんどの場合、すぐ隣のブロックと重複します。すべての読み取りはサンプラーを介して行われます。ホスト プログラムは PyOpenCL (実際には何も変わらないと思います) で、ワーク グループのサイズは 16x16 です。
編集：以下の提案に従ってカーネルの新しいバージョン、作業領域をローカル変数にコピーします：
__kernel __attribute__((reqd_work_group_size(16, 16, 1)))
void test_kernel(
read_only image2d_t src,
write_only image2d_t dest,
const int width,
const int height
)
{
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
int2 pos = (int2)(get_global_id(0), get_global_id(1));
int dx = pos.x % 16;
int dy = pos.y % 16;
__local uint4 local_src[16*16];
__local uint4 local_src2[32*32];
local_src[(pos.y % 16) * 16 + (pos.x % 16)] = read_imageui(src, sampler, pos);
local_src2[(pos.y % 16) * 32 + (pos.x % 16)] = read_imageui(src, sampler, pos);
local_src2[(pos.y % 16) * 32 + (pos.x % 16) + 16] = read_imageui(src, sampler, (int2)(pos.x + 16, pos.y));
local_src2[(pos.y % 16 + 16) * 32 + (pos.x % 16)] = read_imageui(src, sampler, (int2)(pos.x, pos.y + 16));
local_src2[(pos.y % 16 + 16) * 32 + (pos.x % 16) + 16] = read_imageui(src, sampler, (int2)(pos.x + 16, pos.y + 16));
barrier(CLK_LOCAL_MEM_FENCE);
uint4 diff = (uint4)(0, 0, 0, 0);
for (int i=0; i<16; i++)
{
for (int j=0; j<16; j++)
{
diff += local_src[ j*16 + i ] - local_src2[ (j+dy)*32 + i+dx ];
}
}
write_imageui(dest, pos, diff);
}
結果: 出力は正しく、実行時間は56% 遅くなります。(local_src2 ではなく) local_src のみを使用すると、結果は最大 10% 高速になります。
編集: はるかに強力なハードウェアでベンチマークされた、AMD Radeon HD 7850 は 420GFLOPS を取得し、仕様は 1751GFLOPS です。公平を期すために、仕様は乗加算用であり、ここには乗算がないため、予想される値は ~875GFLOPS ですが、理論上のパフォーマンスと比較すると、これはまだかなりずれています。
編集:これを試してみたい人のためにテストを簡単に実行できるように、以下の PyOpenCL のホスト側プログラム:
import pyopencl as cl
import numpy
import numpy.random
from time import time
CL_SOURCE = '''
// kernel goes here
'''
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx, properties=cl.command_queue_properties.PROFILING_ENABLE)
prg = cl.Program(ctx, CL_SOURCE).build()
h, w = 1024, 1024
src = numpy.zeros((h, w, 4), dtype=numpy.uint8)
src[:,:,:] = numpy.random.rand(h, w, 4) * 255
mf = cl.mem_flags
src_buf = cl.image_from_array(ctx, src, 4)
fmt = cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.UNSIGNED_INT8)
dest_buf = cl.Image(ctx, mf.WRITE_ONLY, fmt, shape=(w, h))
# warmup
for n in range(10):
event = prg.test_kernel(queue, (w, h), (16,16), src_buf, dest_buf, numpy.int32(w), numpy.int32(h))
event.wait()
# benchmark
t1 = time()
for n in range(100):
event = prg.test_kernel(queue, (w, h), (16,16), src_buf, dest_buf, numpy.int32(w), numpy.int32(h))
event.wait()
t2 = time()
print "Duration (host): ", (t2-t1)/100
print "Duration (event): ", (event.profile.end-event.profile.start)*1e-9
EDIT : メモリ アクセス パターンを考えると、元の素朴なバージョンはかなり良いかもしれません。ワークグループ内のすべてのワークアイテムを呼び出すとread_imageui(src, sampler, (int2)(pos0.x + i, pos0.y + j))、同じ場所が読み取られます (つまり、これは 1 つの読み取りに過ぎませんか??)。また、呼び出すとread_imageui(src, sampler, (int2)(pos.x + i, pos.y + j))、それらは連続した場所を読み取っています (したがって、読み取りを完全に結合できますか??)。
