OpenCL: ¿es posible invocar otra función desde dentro de un kernel?

Question

estoy siguiendo, junto con un tutorial que se encuentra aquí: http://opencl.codeplex.com/wikipage?title=OpenCL%20Tutorials%20-%201

El kernel que han enumerado es esto, que calcula la suma de dos números y lo almacena en la variable de salida:

__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) Di por ejemplo, que la operación realizada fue mucho más compleja que una sumatoria, algo que garantiza su propia función. Vamos a llamarlo ComplexOp (in1, in2, out). ¿Cómo podría implementar esta función para que vector_add_gpu() pueda llamarla y usarla? ¿Puedes dar un código de ejemplo?

2) Ahora tomemos el ejemplo al extremo, y ahora quiero llamar a una función genérica que opera en los dos números. ¿Cómo podría configurarlo para que se le pueda pasar al núcleo un puntero a esta función y llamarlo según sea necesario?

Answer 1

Puede tener funciones auxiliares para usar en el kernel, consulte OpenCL user defined inline functions. No puede pasar punteros a la función en el kernel.

Answer 2

Sí, es posible. Solo tienes que recordar que OpenCL está basado en C99 con algunas advertencias. Puede crear otras funciones dentro del mismo archivo kernel o en un archivo separado y simplemente incluirlo al principio. No es necesario declarar las funciones auxiliares como en línea, sin embargo, tenga en cuenta que OpenCL alineará las funciones cuando se las llame. Los punteros tampoco están disponibles para usar al llamar a funciones auxiliares.

Ejemplo

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