Compiling a shared library for the MIC?

Compiling a shared library for the MIC?

I'm having a little trouble getting my code up and running on the MIC. My code is C/C++ and is compiled into a series of shared libraries with Python bindings accessible via SWIG. I want to compile some of my kernels in a shared library specifically for the MIC, using the explicit memory model. I'm having trouble offloading a simple for loop inside one of my kernels (a class method inside my solver class). I think that the issue may be with compiling a shared library, part of which is computed on the device and the rest of which runs on the host. I get the following error at runtime:

offload error: cannot find offload entry __offload_entry_MICSolver_cpp_556_ZN9MICSolver15zeroTrackFluxesEvoffload error: process on the device 0 unexpectedly exited with code 1

The class method which I am trying to offload on the device is the following:

void MICSolver::zeroTrackFluxes() {

    int size = 2*_tot_num_tracks*_polar_times_groups;

    #pragma offload target(mic) inout(_boundary_flux:length(size))

    #pragma omp parallel for   

    for (int i=0; i < _tot_num_tracks; i++) {       

        for (int pe2=0; pe2 < 2*_polar_times_groups; pe2++)           

            _boundary_flux[2*i*_polar_times_groups + pe2] = 0.0;                                                      

    }
    return;
}

where MICSolver has "_boundary_flux" (an array allocated on the host), "_tot_num_tracks" and "_polar_times_groups" as class attributes.
Since I'm new to the MIC and can't even offload this simple method, I thought I would see if you have any advice. Thanks for any help you can give!

Best,

Will

4 posts / 0 new
Last post
For more complete information about compiler optimizations, see our Optimization Notice.

The error is typical of an entity lacking __attribute__ ((target (mic))) decoration. I don’t know whether that is missing for your solver class or just the class method. Let me check w/those more expert in C++.

Kevin,

Thanks for the reply - I'm curious to know if I need to decorate the class method and/or the class itself with *__attribute__((target(mic)))*, and if so, if I must do so in both the header and/or the implementation file. I tried this in the header, as well as using the #pragma push/pop decorators around the class method declaration in the header.

I will try a couple perturbations of this later this morning. I can send you my source to show your colleagues if that would be of any help.

Will,

I shared some feedback from Development below but if this is not helpful in resolving the error then we can discuss sharing your code with us.

The Developer said: 

“It may be sufficient to mark the method with the attribute, but if the function is virtual, the class needs to be marked as well.”

While noting being wary of advising based on the code snippets, he said “A routine with a #pragma offload within it is supposed to get marked implicitly target(mic), but if the declarations and definitions are separate, it is possible that all callers may not see that the callee is marked.”

Let me know whether this helps.

Leave a Comment

Please sign in to add a comment. Not a member? Join today