1

OpenCL と image2d_t オブジェクトを使用して画像の畳み込みを高速化しようとしています。出力がすべてゼロの空のイメージであることに気付いたとき、OpenCL カーネルを入力からの基本的な読み取りと出力への書き込みに単純化しました (以下を参照)。少し調整して、画像のいくつかの散らばったピクセルを出力画像に書き込むようにしました。

OpenCL カーネルで read_imageui() を呼び出すまで、イメージが損なわれていないことを確認しました。CommandQueue::enqueueWriteImage() を使用して GPU メモリに画像を書き込み、すぐに CommandQueue::enqueueReadImage() を使用して CPU メモリ内の新しいバッファに読み込みます。この呼び出しの結果は、元の入力画像と一致しました。ただし、カーネルで read_imageui() を使用してピクセルを取得すると、ピクセルの大部分が 0 に設定されます。

C++ ソース:

int height = 112;
int width = 9216;
unsigned int numPixels = height * width;
unsigned int numInputBytes = numPixels * sizeof(uint16_t);
unsigned int numDuplicatedInputBytes = numInputBytes * 4;
unsigned int numOutputBytes = numPixels * sizeof(int32_t);

cl::size_t<3> origin;
origin.push_back(0);
origin.push_back(0);
origin.push_back(0);
cl::size_t<3> region;
region.push_back(width);
region.push_back(height);
region.push_back(1);

std::ifstream imageFile("hri_vis_scan.dat", std::ifstream::binary);
checkErr(imageFile.is_open() ? CL_SUCCESS : -1, "hri_vis_scan.dat");
uint16_t *image = new uint16_t[numPixels];
imageFile.read((char *) image, numInputBytes);
imageFile.close();

// duplicate our single channel image into all 4 channels for Image2D
cl_ushort4 *imageDuplicated = new cl_ushort4[numPixels];
for (int i = 0; i < numPixels; i++)
    for (int j = 0; j < 4; j++)
        imageDuplicated[i].s[j] = image[i];

cl::Buffer imageBufferOut(context, CL_MEM_WRITE_ONLY, numOutputBytes, NULL, &err);
checkErr(err, "Buffer::Buffer()");

cl::ImageFormat inFormat;
inFormat.image_channel_data_type = CL_UNSIGNED_INT16;
inFormat.image_channel_order = CL_RGBA;
cl::Image2D bufferIn(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, inFormat, width, height, 0, imageDuplicated, &err);
checkErr(err, "Image2D::Image2D()");

cl::ImageFormat outFormat;
outFormat.image_channel_data_type = CL_UNSIGNED_INT16;
outFormat.image_channel_order = CL_RGBA;
cl::Image2D bufferOut(context, CL_MEM_WRITE_ONLY, outFormat, width, height, 0, NULL, &err);
checkErr(err, "Image2D::Image2D()");

int32_t *imageResult = new int32_t[numPixels];
memset(imageResult, 0, numOutputBytes);

cl_int4 *imageResultDuplicated = new cl_int4[numPixels];
for (int i = 0; i < numPixels; i++)
    for (int j = 0; j < 4; j++)
        imageResultDuplicated[i].s[j] = 0;

std::ifstream kernelFile("convolutionKernel.cl");
checkErr(kernelFile.is_open() ? CL_SUCCESS : -1, "convolutionKernel.cl");
std::string imageProg(std::istreambuf_iterator<char>(kernelFile), (std::istreambuf_iterator<char>()));
cl::Program::Sources imageSource(1, std::make_pair(imageProg.c_str(), imageProg.length() + 1));
cl::Program imageProgram(context, imageSource);
err = imageProgram.build(devices, "");
checkErr(err, "Program::build()");

cl::Kernel basic(imageProgram, "basic", &err);
checkErr(err, "Kernel::Kernel()");

basic.setArg(0, bufferIn);
basic.setArg(1, bufferOut);
basic.setArg(2, imageBufferOut);

queue.finish();

cl_ushort4 *imageDuplicatedTest = new cl_ushort4[numPixels];
for (int i = 0; i < numPixels; i++)
{
    imageDuplicatedTest[i].s[0] = 0;
    imageDuplicatedTest[i].s[1] = 0;
    imageDuplicatedTest[i].s[2] = 0;
    imageDuplicatedTest[i].s[3] = 0;
}
double gpuTimer = clock();

err = queue.enqueueReadImage(bufferIn, CL_FALSE, origin, region, 0, 0, imageDuplicatedTest, NULL, NULL);
checkErr(err, "CommandQueue::enqueueReadImage()");

// Output from above matches input image

err = queue.enqueueNDRangeKernel(basic, cl::NullRange, cl::NDRange(height, width), cl::NDRange(1, 1), NULL, NULL);
checkErr(err, "CommandQueue::enqueueNDRangeKernel()");

queue.flush();

err = queue.enqueueReadImage(bufferOut, CL_TRUE, origin, region, 0, 0, imageResultDuplicated, NULL, NULL);
checkErr(err, "CommandQueue::enqueueReadImage()");

queue.flush();

err = queue.enqueueReadBuffer(imageBufferOut, CL_TRUE, 0, numOutputBytes, imageResult, NULL, NULL);
checkErr(err, "CommandQueue::enqueueReadBuffer()");

queue.finish();

OpenCL カーネル:

__kernel void basic(__read_only image2d_t input, __write_only image2d_t output, __global int *result)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_TRUE | //Natural coordinates
     CLK_ADDRESS_NONE | //Clamp to zeros
     CLK_FILTER_NEAREST; //Don't interpolate

int2 coord = (get_global_id(1), get_global_id(0));

uint4 pixel = read_imageui(input, smp, coord);
result[coord.s0 + coord.s1 * 9216] = pixel.s0;
write_imageui(output, coord, pixel);
}

カーネルの座標は現在 (x, y) = (幅, 高さ) にマップされています。

入力画像は、1 ピクセルあたり 16 ビットの単一チャネル グレースケール画像です。そのため、OpenCL の Image2D に適合するようにチャネルを複製する必要がありました。畳み込み後の出力はピクセルあたり 32 ビットになるため、numOutputBytes がそれに設定されます。また、幅と高さが奇妙に見えますが、入力画像のサイズは 9216x7824 であるため、最初にコードをテストするためにその一部のみを使用しているため、永遠にかかることはありません。

カーネル内のイメージから読み取った後、グローバル メモリへの書き込みを追加して、問題がイメージの読み取りまたはイメージの書き込みであるかどうかを確認しました。カーネルの実行後、グローバル メモリのこのセクションにもほとんどゼロが含まれます。

どんな助けでも大歓迎です！

4

1 に答える 1

2

read_imageuiのドキュメントには、次のように記載されています。

さらに、整数座標を取る read_imagei および read_imageui 呼び出しは、正規化された座標がCLK_NORMALIZED_COORDS_FALSEに設定され、アドレッシング モードがCLK_ADDRESS_CLAMP_TO_EDGE、CLK_ADDRESS_CLAMP または CLK_ADDRESS_NONE に設定されたサンプラーを使用する必要があります。それ以外の場合、返される値は未定義です。

しかし、CLK_NORMALIZED_COORDS_TRUE を使用してサンプラーを作成しています (ただし、正規化されていない座標を渡しているようです:S ?)。

于 2012-07-18T19:15:49.653 に答える