Indirect function/procedure calls are not supported: seeking workaround

Accelerated Computing HPC Compilers Legacy PGI Compilers
1

Hello,

I am writing to seek some workarounds for the error “Indirect function/procedure calls are not supported”. A simple code is pasted below.

#include <iostream>

class EquationOfStateFR
{
public:
    EquationOfStateFR(){}
    virtual ~EquationOfStateFR(){}
    virtual void createDeviceData(){}
public:
    inline virtual int  computePressure(){return 5;}
};

class EOS_IdealGas : public EquationOfStateFR
{
public:
    inline EOS_IdealGas(){}
    inline ~EOS_IdealGas(){}
    void createDeviceData()
        {
          #pragma acc enter data copyin(this) 
        }
public:
    #pragma acc routine seq
    inline int  computePressure(){return 10;}
};
class solveSystem
{
public:
    solveSystem(){_equationOfState = new EOS_IdealGas;}
    ~solveSystem(){}
    void createDeviceData()
        {
          #pragma acc enter data copyin(this)
          _equationOfState->createDeviceData();
          #pragma acc enter data attach(_equationOfState) 
        }
public:
        #pragma acc routine seq
        void calculation()
        {
          out=_equationOfState->computePressure();
        }
        void print_val()
        {
          #pragma acc update self(out)
          std::cout << "out=" << out << std::endl;
        }
private:
        EOS_IdealGas* _equationOfState;
        int out;
};

int main()
{
        solveSystem obj;
        obj.createDeviceData();
        #pragma acc serial
        {
          obj.calculation();
        }
        obj.print_val();
}

During compilation (with ‘pgc++ -acc -ta=nvidia -Minfo=accel indirect5.C’), the following errors arise:

main:
     59, Accelerator serial kernel generated
         Generating Tesla code
         Generating implicit copy(obj)
EOS_IdealGas::createDeviceData():
     21, Generating enter data copyin(this[:1])
EOS_IdealGas::computePressure():
     24, Generating acc routine seq
         Generating Tesla code
solveSystem::createDeviceData():
     35, Generating enter data copyin(this[:1])
     37, Generating enter data attach(_equationOfState)
PGCC-S-0155-Compiler failed to translate accelerator region (see -Minfo messages): Indirect function/procedure calls are not supported (indirect5.C: 42)
solveSystem::calculation():
     41, Generating acc routine seq
         Generating Tesla code
         42, Accelerator restriction: Indirect function/procedure calls are not supported
PGCC-F-0704-Compilation aborted due to previous errors. (indirect5.C)
PGCC/x86 Linux 18.10-1: compilation aborted

I am aware that currently OpenACC does not support virtual functions in the device code, and it is perhaps why I received this error (could you confirm? It is a bit surprising that even using a derived class pointer is not allowed once virtual functions are present). My intention is therefore to seek your suggestions and recommendations on workarounds. Since we do not want to break the polymorphism by removing the ‘virtual’ qualifier, is it true that our last resort is to create a new function like ‘computePressure_acc’ only in the derived (EOS_IdealGas) class?

Thanks,
Shine

2

I am aware that currently OpenACC does not support virtual functions in the device code, and it is perhaps why I received this error (could you confirm?

Correct, virtual device functions are not yet supported.

It is a bit surprising that even using a derived class pointer is not allowed once virtual functions are present).

Since I’m not a compiler engineer myself, I don’t have a firm grasp on the details. But it’s my understanding that in order to support virtual functions on the device, we would first need to create virtual jump table on the device to support indirect calling. Virtual functions delay binding to a particular function until the function is called, so needsthe jump table to resolve the reference. It’s possible, but difficult to implement. It’s still on our wish list of things to implement, but other features such as Lambda support (done for capture by value) and work on C++17 Parallel STL (in processes) are higher priority.

is it true that our last resort is to create a new function like ‘computePressure_acc’ only in the derived (EOS_IdealGas) class?

Correct. Device functions must be resolved during linking.

Hope this helps,
Mat

3

Thanks, Mat, for your detailed explanations and the heads-up! I am sure the community will benefit a lot from Lambda support and parallel STL. For now I will create dedicated device functions as the workaround.

Thanks,
Shine