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]);
}