OpenMP failing offload of static object if bitfields

[KaruroChori]

Tested in one of the latest commits. This works as intended:

    struct hello{
        struct sub{
            uint32_t uid;
            uint32_t gid ;
            uint32_t idx;
            uint32_t weak ;
        };
        
        #pragma omp declare target
        constexpr static sub A={1,1,1,false};
        #pragma omp end declare target
    };

while this does not:

    struct hello{
        struct sub{
            uint32_t uid:32;
            uint32_t gid:32;
            uint32_t idx:32;
            uint32_t weak:32;
        };
        
        #pragma omp declare target
        constexpr static sub A={1,1,1,false};
        #pragma omp end declare target
    };

Bitfields are intrinsically not working, and the error message is a generic sounding Offloading entry for declare target variable _ZN3sdf5hello1AE is incorrect: the address is invalid.
At the very least, a more explicit description would help but still, I see no reason why bitfields should not be allowed.

Potential workarounds like wrapping it in an union and setting it up through a different field will not work.

[llvmbot]

@llvm/issue-subscribers-openmp

Author: None (KaruroChori)

Tested in one of the latest commits. This works as intended: ``` struct hello{ struct sub{ uint32_t uid; uint32_t gid ; uint32_t idx; uint32_t weak ; };

    #pragma omp declare target
    constexpr static sub A={1,1,1,false};
    #pragma omp end declare target
};

while this does not:

struct hello{
    struct sub{
        uint32_t uid:32;
        uint32_t gid:32;
        uint32_t idx:32;
        uint32_t weak:32;
    };
    
    #pragma omp declare target
    constexpr static sub A={1,1,1,false};
    #pragma omp end declare target
};

Bitfields are intrinsically not working, and the error message is a generic sounding `Offloading entry for declare target variable _ZN3sdf5hello1AE is incorrect: the address is invalid.`
At the very least, a more explicit description would help but still, I see no reason why bitfields should not be allowed.

Potential workarounds like wrapping it in an union and setting it up through a different field will not work.
</details>

[shiltian]

Hmm, wonder what the difference is in terms of IR representation of the two structs.

[KaruroChori]

Honestly, no idea and I have no clue where to start to check that. But I would be mildly surprised if that was the case.
At the very least, I am pretty positive the layout of the final structs is equivalent, or my application would be broken (more broken).
I guess there could be a potential problem with endianess on some combinations of architectures maybe?

[KaruroChori]

More info as I tested this further:

• The issue is not specific for one single target. It will fail for both amd64 and nvptx.
• Bitfields are perfectly usable within target regions. Also copying in and out of target regions seems fine for all scenarios I tested. The problem is specifically when using the declare target clause.

I also tried to compare the llvm bytecode, but I cannot spot any meaningful difference
https://godbolt.org/z/odKE5df1e

[jhuber6]

The error is emitted here https://github.com/llvm/llvm-project/blob/main/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp#L9499. We need an address to be able to emit an offloading entry, because an offloading entry fundamentally is a named symbol in the ELF that the device runtime can look up and copy / read to. A bitfield should have an address, but I'm going to guess that something inside of OpenMP gets confused when it's used. I don't know anything more than that, maybe @alexey-bataev has a guess.

[alexey-bataev]

Bitfields are not addressable by the standard

[KaruroChori]

The bit-aligned fields inside a structure surely are not, but I would assume there should be no problem addressing full instances of structs, even if they have bitfields inside.
Is it failing because it wants each field to be addressable?

[KaruroChori]

Just a quick addendum.
This issue is quite annoying as it will also impact objects which are nested inside other objects. If one is vendoring all code and libraries themselves it can be worked around, but in general this issue breaks compatibility with arbitrary libraries.

As for the workaround:

struct bitfield{
    uint32_t a:16;
    uint32_t b:16;
};

template<typename T>
struct wrapper{
    uint8_t data[sizeof(T)];
    operator T&(){return *(T*)this;}
    wrapper(const T& r){*(T*)this=r;}
    T* operator->(){return (T*)this;}
};

#pragma omp declare target
static wrapper<bitfield> a = bitfield{4,6}; //This is ok
static bitfield b = {5,6};                  //This one fails
#pragma omp end declare target

int main(){
    #pragma omp target
    {   
        printf("%d %d",a->a, a->b);
        printf("%d %d",b.a, b.b);
    }
    return 0;
}

It is possible to introduce a template wrapper. "Surprisingly" this works, as annoying as the extra syntax needed is.