CudafyCompileException OpenCL

Mar 6, 2014 at 6:41 AM
I am trying to compile the below c# code for OpenCL but I get a CudafyCompileException. I post the c# code, the c source code as translated by Cudafy and the compiler output. The same code compiles well for Cuda and Emulator. Any help appreciated!

C# Code

using Cudafy;
using Cudafy.Host;
using Cudafy.Translator;
using Cudafy.Types;
using System;
using System.Collections.Generic;
using System.Numerics;
using System.Linq;

namespace GPU_Functions
{
public static class SampleOpenCL
{
    [Cudafy]
    public struct ComplexFloat
    {
        public float x;
        public float y;
    }

    [Cudafy]
    public static ComplexFloat Add(ComplexFloat x, ComplexFloat y)
    {
        x.x = x.x + y.x;
        x.y = x.y + y.y;
        return x;
    }

    public static ComplexFloat[,] Calculate(List<float[,]> dataList, float[] frequencies, Cudafy.eGPUType type)
    {

       CudafyTranslator.Language = type == eGPUType.OpenCL ? eLanguage.OpenCL : eLanguage.Cuda;
        CudafyModule km = CudafyTranslator.Cudafy();
        var _gpu = CudafyHost.GetDevice(type);

        _gpu.LoadModule(km);

        ComplexFloat[,] totalResults = new ComplexFloat[dataList.Sum(x => x.GetLength(0)), frequencies.Length];
        int index = 0;
        foreach (var data in dataList)
        {
            var pathsCount = data.GetLength(0);
            ComplexFloat[,] results = new ComplexFloat[pathsCount, frequencies.Length];

            float[] dev_frequencies = _gpu.CopyToDevice(frequencies);
            float[,] dev_data = _gpu.CopyToDevice(data);

            ComplexFloat[,] dev_results = _gpu.CopyToDevice<ComplexFloat>(results);
            _gpu.Launch(pathsCount, frequencies.Length).CalculatePathsSW(dev_frequencies, dev_data, dev_results);
            _gpu.CopyFromDevice(dev_results, results);

            for (int i = 0; i < pathsCount; i++)
            {
                for (int j = 0; j < frequencies.Length; j++)
                    totalResults[index + i, j] = results[i, j];// new ComplexFloat(results[i, j].x, results[i, j].y);
            }
            index = index + pathsCount;
        }
        _gpu.FreeAll();
        return totalResults;
     }


    [Cudafy]
    public static void CalculateFrequencies(GThread thread, ComplexFloat[,] results, ComplexFloat[] frequencyResults)
    {
        int bid = thread.blockIdx.x;
        int pathsCount = results.GetLength(0);
        for (int i = 0; i < pathsCount; i++)
        {

            frequencyResults[bid] = Add(frequencyResults[bid], results[i, bid]);
        }
    }

    [Cudafy]
    public static void CalculatePaths(GThread thread, float[] frequencies, float[,] data, ComplexFloat[,] result)
    {
        int bid = thread.blockIdx.x;
        int tid = thread.threadIdx.x;
        int distanceCol = data.GetLength(1) - 1;
        int speedOfSoundCol = data.GetLength(1) - 2;
        int airDensityCol = data.GetLength(1) - 3;
        float distance = data[bid, distanceCol];
        float airDensity = data[bid, airDensityCol];
        float speedOfSound = data[bid, speedOfSoundCol];
        float frequency = frequencies[tid];
        //wavenumber calculation
        float k = (2 * Cudafy.GMath.PI / speedOfSound * frequency);
        //distance attenuation calculation
        ComplexFloat pressure = new ComplexFloat(); // new ComplexFloat(Cudafy.GMath.Cos(k * distance), Cudafy.GMath.Sin(k * distance));
        result[bid, tid] = pressure;
    }     
}
}

Cudafy OpenCL Source Code

if defined(cl_khr_fp64)

pragma OPENCL EXTENSION cl_khr_fp64: enable

elif defined(cl_amd_fp64)

pragma OPENCL EXTENSION cl_amd_fp64: enable

endif

struct SampleOpenCLComplexFloat
{
float x;
float y;
};


// GPU_Functions.SampleOpenCL
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y);
// GPU_Functions.SampleOpenCL
__kernel void CalculateFrequencies(global struct SampleOpenCLComplexFloat* results, int resultsLen0, int resultsLen1, global struct SampleOpenCLComplexFloat* frequencyResults, int frequencyResultsLen0);
// GPU_Functions.SampleOpenCL
__kernel void CalculatePaths(global float* frequencies, int frequenciesLen0, global float* data, int dataLen0, int dataLen1, global struct SampleOpenCLComplexFloat* result, int resultLen0, int resultLen1);

// GPU_Functions.SampleOpenCL
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y)
{
x.x += y.x;
x.y += y.y;
return x;
}
// GPU_Functions.SampleOpenCL
__kernel void CalculateFrequencies(global struct SampleOpenCLComplexFloat* results, int resultsLen0, int resultsLen1, global struct SampleOpenCLComplexFloat* frequencyResults, int frequencyResultsLen0)
{
int x = get_group_id(0);
int length = resultsLen0;
for (int i = 0; i < length; i++)
{
    frequencyResults[(x)] = Add(frequencyResults[(x)], results[(i) * resultsLen1 + ( x)]);
}
}
// GPU_Functions.SampleOpenCL
__kernel void CalculatePaths(global float* frequencies, int frequenciesLen0, global float* data, int dataLen0, int dataLen1, global struct SampleOpenCLComplexFloat* result, int resultLen0, int resultLen1)
{
int x = get_group_id(0);
int x2 = get_local_id(0);
int num = dataLen1 - 1;
int num2 = dataLen1 - 2;
int num3 = dataLen1 - 3;
float num4 = data[(x) * dataLen1 + ( num)];
float num5 = data[(x) * dataLen1 + ( num3)];
float num6 = data[(x) * dataLen1 + ( num2)];
float num7 = frequencies[(x2)];
float num8 = 6.283185f / num6 * num7;
SampleOpenCLComplexFloat complexFloat = SampleOpenCLComplexFloat();
result[(x) * resultLen1 + ( x2)] = complexFloat;
}

Compiler output

:14:3: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y);
^
struct
:14:32: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y);
                           ^
                           struct 
:14:60: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y);
                                                       ^
                                                       struct 
:21:3: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y)
^
struct
:21:32: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y)
                           ^
                           struct 
:21:60: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
SampleOpenCLComplexFloat Add(SampleOpenCLComplexFloat x, SampleOpenCLComplexFloat y)
                                                       ^
                                                       struct 
:50:2: error: must use 'struct' tag to refer to type 'SampleOpenCLComplexFloat'
    SampleOpenCLComplexFloat complexFloat = SampleOpenCLComplexFloat();
    ^
    struct 
:50:42: error: implicit declaration of function 'SampleOpenCLComplexFloat' is invalid in OpenCL
    SampleOpenCLComplexFloat complexFloat = SampleOpenCLComplexFloat();
                                            ^
:50:27: error: initializing 'struct SampleOpenCLComplexFloat' with an expression of incompatible type 'int'
    SampleOpenCLComplexFloat complexFloat = SampleOpenCLComplexFloat();
                             ^              ~~~~~~~~~~~~~~~~~~~~~~~~~~
error: front end compiler failed build.
Coordinator
Mar 6, 2014 at 8:30 AM
You need to give the translator some hints in the format of Cudafy____ attributes on the function parameters. Some steps such as the last line above plain do not work correctly. The whole OpenCL support is rather limited compared to CUDA.
Mar 7, 2014 at 8:04 AM
Dear NickKopp,

Thanks for the reply but I am afraid I did not understand what I am supposed to do. Can you elaborate a bit more?