Недавно я портировал мой volumeraycaster из OpenGL в OpenCL, что снизило производительность raycaster примерно на 90 процентов. Я отслеживал снижение производительности с помощью функций обработки изображений OpenCL, которые намного медленнее, чем соответствующие функции текстурирования OpenGL. Удалив функции репликации изображений и функции выборки текстуры, обе версии raycaster имели примерно такую же скорость. Чтобы легко сканировать функции на разных аппаратных средствах и исключать некоторые глупые ошибки в остальном коде моего RT, я написал небольшой тест, который сравнивает скорость выборки OpenCL с частотой выборки OpenGL и тестировал ее на разных машинах, но В OpenCL все еще было около 10% производительности OpenGL.
Тест OpenCL HostCode (по крайней мере, самая важная его часть):
void OGLWidget::OCLImageSampleTest()
{
try
{
int size=8;
float Values[4*size*size*size];
cl::Kernel kernel=cl::Kernel(program,"ImageSampleTest",NULL);
cl::ImageFormat FormatA(CL_RGBA,CL_FLOAT);
cl::Image3D CLImage(CLcontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR ,FormatA,size,size,size,0,0,Values,NULL);
cl::ImageFormat FormatB(CL_RGBA,CL_UNSIGNED_INT8);
cl::Image2D TempImage(CLcontext, CL_MEM_WRITE_ONLY,FormatB,1024,1024,0,NULL,NULL );
kernel.setArg(0, CLImage);
kernel.setArg(1, TempImage);
cl::Sampler Samp;
Samp() = clCreateSampler( CLcontext(), CL_TRUE, CL_ADDRESS_REPEAT, CL_FILTER_LINEAR, NULL);
kernel.setArg(2, Samp);
QTime BenchmarkTimer=QTime();
BenchmarkTimer.start();
cl::KernelFunctor func = kernel.bind(queue, cl::NDRange(1024,1024), cl::NDRange(32,32));
func().wait();
int Duration = BenchmarkTimer.elapsed();
printf("OCLImageSampleTest: %d ms \n", Duration);
}
catch (cl::Error& err)
{
std::cerr << "An OpenCL error occured, " << err.what()
<< "\nError num of " << err.err() << "\n";
return;
}
}
OpenCL Kernel:
void kernel ImageSampleTest( read_only image3d_t CoordTexture, write_only image2d_t FrameBuffer, sampler_t smp)
{
int Screenx = get_global_id(0);
int Screeny = get_global_id(1);
int2 PositionOnScreen=(int2)(Screenx,Screeny) ;
float4 Testvec=(float4)(1,1,1,1);
for(int i=0; i< 2000; i++)
{
Testvec+= read_imagef(CoordTexture,smp, (float4)(0+0.00000001*i,0,0,0)); // i makes sure that the compiler doesn't unroll the loop
}
uint4 ToInt=(uint4)( (uint) (Testvec.x), (uint) (Testvec.y) ,(uint)(Testvec.z),1);
write_imageui ( FrameBuffer, PositionOnScreen, ToInt );
}
OpenGL FragmentShader для полноэкранного квадроцикла, который имеет такое же количество фрагментов, что и ядро OpenCL, имеет рабочие элементы:
#version 150
uniform sampler3D Tex;
out vec4 FragColor;
void main()
{
FragColor=vec4(0,0,0,0);
for(int i=0; i<2000; i++)
{
FragColor+= texture(Tex,vec3(0+0.00000001*i,0,0),0);
}
}
Кроме того, я уже пробовал следующее повысить производительность:
-изменяет размер рабочей группы: нет увеличения производительности
-Different Hardware: 280 GTX, 580 GTX, некоторые карты Fermi Tessla, все они имели ту же ужасную производительность в OpenCL vs OpenGL
- Различные форматы текстур (байты вместо поплавков), разные шаблоны доступа и разные размеры текстур: без увеличения
-Использование буфера вместо изображения для данных и самонаписанной трилинейной интерполяционной функции для выборки в ядре CL: увеличение производительности OpenCL примерно на 100%
-Использование 2D-изображения//текстуры вместо 3D-изображения//texture: это увеличило производительность OpenCL на 100%, хотя производительность OpenGL вообще не изменилась.
-Использование "ближайшего" вместо "линейной" интерполяции: без изменения производительности
Это оставило меня интересно: Я сделал очень тупую ошибку, которая снижает производительность OpenCL? Почему производительность выборки OpenCL настолько низкая, хотя она должна использовать те же аппаратные средства, что и OpenGL? Почему моя сложная трилинейная функция интерполяции быстрее, чем ее аппаратная реализация? Как увеличить производительность выборки в OpenCL, чтобы я мог иметь ту же скорость, что и в OpenGL?