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 用户定义的内联函数。您不能将函数指针传递到内核中。