Hi everyone. I am porting a CUDA kernel to OpenCL kernel to execute it on a FPGA (Stratix V). The code looks very similar so it shouldn't be a problem except that this FPGA board supports only OpenCL 1.0. So I have some questions regarding the code I have to port. 1) With CUDA you can call different functions asynchronously and allocate and store the data on the GPU regardless of the kernel you are executing (as far as I understood). Is this possible with OpenCL? 2) This algorithm is partially sequential and parallel. How can I achieve the same behavior without loosing performances with OpenCL? Can I use more kernels? And if yes how does that work? 3) Any suggestion about how should I approach this? Thank you very much for those who will be able to help me.

Ok after a long research now I have a good idea about how to port a CUDA code to Opencl. The only thing is now bothering me is this: Assume I have kernels and they call __device__ functions in CUDA which do not exist in OpenCL but they are simple functions called by the kernel and these functions lunches other kernels with a particular global dimension and local dimension, how can I translate this easily into OpenCL? I know that with OpenCL 1.0 I cannot launch kernels from kernels so what's the best way to go about this? Thanks..

1) You can create multiple queues and have multiple kernels or buffer copy calls running in parallel, and then synchronize them using events. Note that global memory consistency is only guaranteed at the end of kernel execution and hence, if you try to share a read_write buffer between multiple queues, you will get undefined behavior, unless you synchronize the calls using events. 2) You can have as many kernels as you want in each queue. Kernel execution calls are non-blocking, and kernels will be queued on the device and executed in-order. You can put the two different parts of your code in two different kernels and run them in-order. 3) I would remove all GPU-based optimizations from the code and convert the kernel to a sequential C code and compile it as a single work-item kernel. You probably need to spend a good deal of time reading Khronos's OpenCL documents and also Altera's OpenCL documents (getting started guide, programming guide, best practices guide). Don't concern yourself with what version of OpenCL is supported on the FPGA, you should not need to use the features that only exist in newer versions just yet.

I am struggling to understand how to covert a NDRange for OpenCL from CUDA. Say this is my kernel launch for CUDA: initMatrix << < _h_cols, _h_rows >> > (var1, va2, var3, var4); Cols and Rows are just integer numbers and the matrix has been made flat so no 2D. How can I convert this to a OpenCL NDRange? Thanks..

If your issue is with general OpenCL usage, and not problems specific to using OpenCL on FPGAs, I recommend searching online or consulting with stackoverflow. You could take a look at the following pages for starters: https://github.com/vtsynergy/cu2cl https://www.sharcnet.ca/help/index.php/porting_cuda_to_opencl

The problem I am encountering is that the number of columns and rows of my matrix are not multiples or dividers of each other. Can someone help me with this?

Help with porting a CUDA code to OpenCL | Altera Community

19 Replies

Altera_Forum
Honored Contributor
9 years ago
Ok after a long research now I have a good idea about how to port a CUDA code to Opencl. The only thing is now bothering me is this:

Assume I have kernels and they call __device__ functions in CUDA which do not exist in OpenCL but they are simple functions called by the kernel and these functions lunches other kernels with a particular global dimension and local dimension, how can I translate this easily into OpenCL? I know that with OpenCL 1.0 I cannot launch kernels from kernels so what's the best way to go about this?

Thanks..
Altera_Forum
Honored Contributor
9 years ago
1) You can create multiple queues and have multiple kernels or buffer copy calls running in parallel, and then synchronize them using events. Note that global memory consistency is only guaranteed at the end of kernel execution and hence, if you try to share a read_write buffer between multiple queues, you will get undefined behavior, unless you synchronize the calls using events.

2) You can have as many kernels as you want in each queue. Kernel execution calls are non-blocking, and kernels will be queued on the device and executed in-order. You can put the two different parts of your code in two different kernels and run them in-order.

3) I would remove all GPU-based optimizations from the code and convert the kernel to a sequential C code and compile it as a single work-item kernel.

You probably need to spend a good deal of time reading Khronos's OpenCL documents and also Altera's OpenCL documents (getting started guide, programming guide, best practices guide). Don't concern yourself with what version of OpenCL is supported on the FPGA, you should not need to use the features that only exist in newer versions just yet.
Altera_Forum
Honored Contributor
9 years ago
I am struggling to understand how to covert a NDRange for OpenCL from CUDA.

Say this is my kernel launch for CUDA: initMatrix << < _h_cols, _h_rows >> > (var1, va2, var3, var4);

Cols and Rows are just integer numbers and the matrix has been made flat so no 2D.
How can I convert this to a OpenCL NDRange?

Thanks..
Altera_Forum
Honored Contributor
9 years ago
If your issue is with general OpenCL usage, and not problems specific to using OpenCL on FPGAs, I recommend searching online or consulting with stackoverflow. You could take a look at the following pages for starters:

https://github.com/vtsynergy/cu2cl

https://www.sharcnet.ca/help/index.php/porting_cuda_to_opencl
Altera_Forum
Honored Contributor
9 years ago
The problem I am encountering is that the number of columns and rows of my matrix are not multiples or dividers of each other. Can someone help me with this?

Altera_Forum

Honored Contributor

9 years ago

Ok so I am getting far with this project and porting is getting easier. But now I am experiencing this problem when I run this kernel:

__kernel void computeRouteStrength(__global double* restrict routeStrength,__global int* restrict order ,__global double* restrict routePheromones,__global Route* restrict possibleRoutesPerOrderMatrix, 
                                    __global int* restrict cols, __global double* restrict beta, __global int* restrict dailyOrderLimit, __global int* restrict productionCapacities)
{
    int i = get_global_id(0);
    int route = (*order * *cols) + i;
    
    double strength = -1;
    if (possibleRoutesPerOrderMatrix.order != -1)
    {
        double pheromone = routePheromones;
        double heuristicInformation = 1/possibleRoutesPerOrderMatrix.heuristicInformation;
        if (canBeShipped(&possibleRoutesPerOrderMatrix, dailyOrderLimit, productionCapacities))
        {
            strength = pheromone * heuristicInformation;
        }
    }
    routeStrength = strength;
}

So basically if I remove the calculation between the if statement the kernel doesn't fail. If I do put any calculation there then it fails saying: INVALID COMMAND QUEUE.

What could be the problem? The index should not be a problem also the arrays have been flattened (from 2d to 1d) and I am using them in other kernels.

Thanks..

Altera_Forum
Honored Contributor
9 years ago
Unless you are converting the OpenCL error number incorrectly, "INVALID COMMAND QUEUE" has nothing to do with what is in your kernel; it is because you have some issue in your host code. Find the OpenCL function in the host that is throwing this error and fix it.
Altera_Forum
Honored Contributor
9 years ago
--- Quote Start ---
Unless you are converting the OpenCL error number incorrectly, "INVALID COMMAND QUEUE" has nothing to do with what is in your kernel; it is because you have some issue in your host code. Find the OpenCL function in the host that is throwing this error and fix it.
--- Quote End ---

I have used clerrors.h from the tests project of the fpga board to decode the error. Also if I run the kernel with some data it doesn't fail. If I do with other data it does fail. I am not sure why. Data doesn't look to be corrupted.
Altera_Forum
Honored Contributor
9 years ago
Apparently that kernel does run if make sequential. As soon as I run it as a parallel execution it fails. Yes I have made changes in the code for each type of execution.
So what I am thinking is a race condition when writing to the same array. Is that correct/possible?

I am trying to use atomic functions to see if I can solve this problem but when I compile I get "note: candidate function has been explicitly made unavailable". What does it mean?
I am using double variables in my code.

Thanks
Altera_Forum
Honored Contributor
9 years ago
As I mentioned earlier:

--- Quote Start ---
Note that global memory consistency is only guaranteed at the end of kernel execution and hence, if you try to share a read_write buffer between multiple queues, you will get undefined behavior, unless you synchronize the calls using events.
--- Quote End ---

So, yes, do not do this unless you are using events to synchronize kernel calls from the host.

Regarding the "note: candidate function has been explicitly made unavailable" message, Altera's Programming Guide says:

The SDK does not support 64-bit atomic functions described in Section 9.7 of the OpenCL Specification version 1.0.

That is probably the source of your problem.

Forum Discussion

Help with porting a CUDA code to OpenCL

19 Replies

Recent Discussions

IP-CXLTYP3DDR-NC

Cyclone 10LP 080 Timing Closure

How to make Questa show delays through gates?

Compiling quartus design of PCIe root port example failing in WSL2

Fatal Error during Place and Route