OpenCL - có thể gọi hàm khác từ bên trong hạt nhân không?

Question

tôi sau cùng với một hướng dẫn nằm ở đây: http://opencl.codeplex.com/wikipage?title=OpenCL%20Tutorials%20-%201

Kernel họ đã liệt kê là thế này, trong đó tính tổng của hai số và lưu trữ nó trong biến đầu ra:

__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) Say ví dụ như các hoạt động thực hiện phức tạp hơn nhiều so với một tổng kết - một cái gì đó đảm bảo chức năng riêng của nó. Hãy gọi nó là ComplexOp (in1, in2, out). Làm thế nào tôi sẽ đi về thực hiện chức năng này như vậy mà vector_add_gpu() có thể gọi và sử dụng nó? Bạn có thể cung cấp mã ví dụ không?

2) Bây giờ chúng ta hãy lấy ví dụ đến cùng cực, và bây giờ tôi muốn gọi một hàm chung hoạt động trên hai số. Làm thế nào tôi sẽ thiết lập nó để hạt nhân có thể được thông qua một con trỏ đến chức năng này và gọi nó là cần thiết?

Answer 1

Bạn có thể có các chức năng phụ trợ để sử dụng trong hạt nhân, xem OpenCL user defined inline functions. Bạn không thể chuyển con trỏ hàm vào kernel.

Answer 2

Có thể thực hiện được. Bạn chỉ cần nhớ rằng OpenCL được dựa trên C99 với một số cảnh báo. Bạn có thể tạo các hàm khác bên trong cùng một tệp kernel hoặc trong một tệp riêng biệt và chỉ bao gồm nó ngay từ đầu. Các hàm phụ trợ không cần phải được khai báo là inline, tuy nhiên, hãy nhớ rằng OpenCL sẽ inline các hàm khi được gọi. Con trỏ cũng không có sẵn để sử dụng khi gọi các chức năng phụ trợ.

Ví dụ

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