以GPU(OpenGL)为目标的Halide-基准测试并使用Halide RuntimeOpenGL.h
Halide with GPU (OpenGL) as Target - benchmarking and using HalideRuntimeOpenGL.h
我是Halide的新手。我一直在玩教程,以了解语言。现在,我正在编写一个小的演示应用程序,从OSX上的命令行运行。
我的目标是对图像执行逐像素操作,在GPU上调度并测量性能。我已经尝试了一些我想在这里分享的事情,并对下一步有一些问题。
第一种方法
我在GPU上调度了算法,目标是OpenGL,但由于我无法访问GPU内存来写入文件,在Halide例程中,我通过创建类似于Halide repo 中的glsl示例应用程序的Func cpu_out
将输出复制到CPU
pixel_operation_cpu_out.cpp
#include "Halide.h"
#include <stdio.h>
using namespace Halide;
const int _number_of_channels = 4;
int main(int argc, char** argv)
{
ImageParam input8(UInt(8), 3);
input8
.set_stride(0, _number_of_channels) // stride in dimension 0 (x) is three
.set_stride(2, 1); // stride in dimension 2 (c) is one
Var x("x"), y("y"), c("c");
// algorithm
Func input;
input(x, y, c) = cast<float>(input8(clamp(x, input8.left(), input8.right()),
clamp(y, input8.top(), input8.bottom()),
clamp(c, 0, _number_of_channels))) / 255.0f;
Func pixel_operation;
// calculate the corresponding value for input(x, y, c) after doing a
// pixel-wise operation on each each pixel. This gives us pixel_operation(x, y, c).
// This operation is not location dependent, eg: brighten
Func out;
out(x, y, c) = cast<uint8_t>(pixel_operation(x, y, c) * 255.0f + 0.5f);
out.output_buffer()
.set_stride(0, _number_of_channels)
.set_stride(2, 1);
input8.set_bounds(2, 0, _number_of_channels); // Dimension 2 (c) starts at 0 and has extent _number_of_channels.
out.output_buffer().set_bounds(2, 0, _number_of_channels);
// schedule
out.compute_root();
out.reorder(c, x, y)
.bound(c, 0, _number_of_channels)
.unroll(c);
// Schedule for GLSL
out.glsl(x, y, c);
Target target = get_target_from_environment();
target.set_feature(Target::OpenGL);
// create a cpu_out Func to copy over the data in Func out from GPU to CPU
std::vector<Argument> args = {input8};
Func cpu_out;
cpu_out(x, y, c) = out(x, y, c);
cpu_out.output_buffer()
.set_stride(0, _number_of_channels)
.set_stride(2, 1);
cpu_out.output_buffer().set_bounds(2, 0, _number_of_channels);
cpu_out.compile_to_file("pixel_operation_cpu_out", args, target);
return 0;
}
由于我编译了这个AOT,所以我在main()
中为它调用了一个函数。main()
驻留在另一个文件中。
main_file.cpp
注意:此处使用的Image
类与此Halide示例应用程序中的类相同
int main()
{
char *encodeded_jpeg_input_buffer = read_from_jpeg_file("input_image.jpg");
unsigned char *pixelsRGBA = decompress_jpeg(encoded_jpeg_input_buffer);
Image input(width, height, channels, sizeof(uint8_t), Image::Interleaved);
Image output(width, height, channels, sizeof(uint8_t), Image::Interleaved);
input.buf.host = &pixelsRGBA[0];
unsigned char *outputPixelsRGBA = (unsigned char *)malloc(sizeof(unsigned char) * width * height * channels);
output.buf.host = &outputPixelsRGBA[0];
double best = benchmark(100, 10, [&]() {
pixel_operation_cpu_out(&input.buf, &output.buf);
});
char* encoded_jpeg_output_buffer = compress_jpeg(output.buf.host);
write_to_jpeg_file("output_image.jpg", encoded_jpeg_output_buffer);
}
这工作得很好,给了我期望的输出。据我所知,cpu_out
使out
中的值在CPU内存中可用,这就是为什么我能够通过访问main_file.cpp
中的output.buf.host
来访问这些值
第二种方法:
我尝试的第二件事是不通过创建Func cpu_out
从Halide时间表中的设备复制到主机,而是使用main_file.cpp
中的copy_to_host
函数。
pixel_operation_gpu_out.cpp
#include "Halide.h"
#include <stdio.h>
using namespace Halide;
const int _number_of_channels = 4;
int main(int argc, char** argv)
{
ImageParam input8(UInt(8), 3);
input8
.set_stride(0, _number_of_channels) // stride in dimension 0 (x) is three
.set_stride(2, 1); // stride in dimension 2 (c) is one
Var x("x"), y("y"), c("c");
// algorithm
Func input;
input(x, y, c) = cast<float>(input8(clamp(x, input8.left(), input8.right()),
clamp(y, input8.top(), input8.bottom()),
clamp(c, 0, _number_of_channels))) / 255.0f;
Func pixel_operation;
// calculate the corresponding value for input(x, y, c) after doing a
// pixel-wise operation on each each pixel. This gives us pixel_operation(x, y, c).
// This operation is not location dependent, eg: brighten
Func out;
out(x, y, c) = cast<uint8_t>(pixel_operation(x, y, c) * 255.0f + 0.5f);
out.output_buffer()
.set_stride(0, _number_of_channels)
.set_stride(2, 1);
input8.set_bounds(2, 0, _number_of_channels); // Dimension 2 (c) starts at 0 and has extent _number_of_channels.
out.output_buffer().set_bounds(2, 0, _number_of_channels);
// schedule
out.compute_root();
out.reorder(c, x, y)
.bound(c, 0, _number_of_channels)
.unroll(c);
// Schedule for GLSL
out.glsl(x, y, c);
Target target = get_target_from_environment();
target.set_feature(Target::OpenGL);
std::vector<Argument> args = {input8};
out.compile_to_file("pixel_operation_gpu_out", args, target);
return 0;
}
main_file.cpp
#include "pixel_operation_gpu_out.h"
#include "runtime/HalideRuntime.h"
int main()
{
char *encodeded_jpeg_input_buffer = read_from_jpeg_file("input_image.jpg");
unsigned char *pixelsRGBA = decompress_jpeg(encoded_jpeg_input_buffer);
Image input(width, height, channels, sizeof(uint8_t), Image::Interleaved);
Image output(width, height, channels, sizeof(uint8_t), Image::Interleaved);
input.buf.host = &pixelsRGBA[0];
unsigned char *outputPixelsRGBA = (unsigned char *)malloc(sizeof(unsigned char) * width * height * channels);
output.buf.host = &outputPixelsRGBA[0];
double best = benchmark(100, 10, [&]() {
pixel_operation_gpu_out(&input.buf, &output.buf);
});
int status = halide_copy_to_host(NULL, &output.buf);
char* encoded_jpeg_output_buffer = compress_jpeg(output.buf.host);
write_to_jpeg_file("output_image.jpg", encoded_jpeg_output_buffer);
return 0;
}
所以,现在,我认为正在发生的是,pixel_operation_gpu_out
将output.buf
保持在GPU上,当我执行copy_to_host
时,这就是我将内存复制到CPU的时候。这个程序也给了我预期的输出。
问题:
第二种方法比第一种方法慢得多。不过,慢速部分不在基准部分。例如,对于第一种方法,我得到17毫秒作为4k图像的基准时间。对于相同的图像,在第二种方法中,我得到基准时间为22us,copy_to_host
所用的时间为10s。我不确定这种行为是否是预期的,因为方法1和方法2本质上都在做相同的事情。
我尝试的下一件事是使用[HalideRuntimeOpenGL.h][3]
并将纹理链接到输入和输出缓冲区,以便能够从main_file.cpp
直接绘制到OpenGL上下文,而不是保存到jpeg文件。然而,我找不到任何例子来说明如何使用HalideRuntimeOpenGL.h
中的函数,无论我自己尝试了什么,都会给我带来运行时错误,我无法解决这些错误。如果有人能为我提供任何资源,那就太好了。
此外,我们也欢迎对以上代码的任何反馈。我知道这很有效,而且正在做我想做的事,但这可能是完全错误的做法,我也不知道会有什么更好的方法。
10拷贝回内存的原因很可能是因为GPU API已将所有内核调用排入队列,然后在调用halide _copy_to_host时等待它们完成。在运行所有计算调用后,您可以在基准计时内调用halide_device_sync来处理获取循环内的计算时间而不需要回拷贝时间。
我无法从代码中判断内核从该代码中运行了多少次。(我的猜测是100,但可能是那些用于基准测试的参数设置了某种参数化,它试图尽可能多次地运行它,以获得显著性。如果是这样,那就是一个问题,因为排队调用非常快,但计算当然是异步的。如果是这种情况,你可以做一些事情,比如排队10个调用,然后调用halide_device_sync并使用数字"10"以真实了解需要多长时间。)
- 使用rdtsc进行基准测试的缺点是什么
- 对 'std::thread::_M_start_thread CMake 的未定义引用进行基准测试
- 更高效地在微控制器上对C++进行基准测试
- _mm256_load_ps调试模式下导致谷歌/基准测试的分段错误
- 二叉树基准测试结果
- 如何使用谷歌基准测试对自定义界面进行基准测试
- 谷歌基准测试,如何只调用一次代码?
- 使用 std::chrono::steady_clock 对线程/异步中的代码进行基准测试
- 谷歌基准测试结果中显示的时间没有意义
- 使用 Google 基准测试时返回值会发生什么情况?
- 如何在Qt测试框架中对信号进行基准测试?
- C/C++memcpu基准测试:测量CPU和墙时间
- 如何将参数传递给Google基准测试程序
- 如何对CUDA项目进行基准测试
- 为什么这个简单的 C++ SIMD 基准测试在使用 SIMD 指令时运行速度较慢?
- 多部分基准测试的权重是多少?
- 简单的 for() 循环基准测试在任何循环绑定下花费相同的时间
- 禁用优化后,quick-bench.com 基准测试要快得多
- Winsock本地客户服务器基准测试
- 以GPU(OpenGL)为目标的Halide-基准测试并使用Halide RuntimeOpenGL.h