OpenCL - возможно ли вызвать другую функцию из ядра?

Question

Я следую наряду с учебником, расположенным здесь: http://opencl.codeplex.com/wikipage?title=OpenCL%20Tutorials%20-%201

Ядро, которое они перечислили, это, которое вычисляет сумму двух чисел и сохраняет ее в выходной переменной:

__kernel void vector_add_gpu (__global const float* src_a,
                     __global const float* src_b,
                     __global float* res,
           const int num)
{
   /* get_global_id(0) returns the ID of the thread in execution.
   As many threads are launched at the same time, executing the same kernel,
   each one will receive a different ID, and consequently perform a different computation.*/
   const int idx = get_global_id(0);

   /* Now each work-item asks itself: "is my ID inside the vector's range?"
   If the answer is YES, the work-item performs the corresponding computation*/
   if (idx < num)
      res[idx] = src_a[idx] + src_b[idx];
}

1) Скажем, например, что выполненная операция была гораздо более сложной, чем суммирование - что-то, что требует своей собственной функции.Давайте назовем это ComplexOp (in1, in2, out).Как мне реализовать эту функцию, чтобы vector_add_gpu () могла вызывать и использовать ее?Можете ли вы привести пример кода?

2) Теперь давайте рассмотрим пример до крайности, и теперь я хочу вызвать обобщенную функцию, которая работает с двумя числами.Как бы настроить его так, чтобы ядру можно было передать указатель на эту функцию и вызывать ее по мере необходимости?

Answer 1

Да, это возможно.Вы просто должны помнить, что OpenCL основан на C99 с некоторыми оговорками.Вы можете создавать другие функции либо внутри того же файла ядра, либо в отдельном файле и просто включить его в начале.Вспомогательные функции не нужно объявлять как встроенные, однако имейте в виду, что OpenCL встроит функции при вызове.Указатели также недоступны для использования при вызове вспомогательных функций.

Пример

float4 hit(float4 ray_p0, float4 ray_p1, float4 tri_v1, float4 tri_v2, float4 tri_v3)
{
//logic to detect if the ray intersects a triangle
}

__kernel void detection(__global float4* trilist, float4 ray_p0, float4 ray_p1)
{
int gid = get_global_id(0);
float4 hitlocation = hit(ray_p0, ray_p1, trilist[3*gid], trilist[3*gid+1], trilist[3*gid+2]);
}

Answer 2

Вы можете иметь вспомогательные функции для использования в ядре, см. Определяемые пользователем встроенные функции OpenCL .Вы не можете передавать указатели функций в ядро.