Bug with nested device data environments implementation

[Alexandros_P_]

According to OpenMP 4.0 specification page 177 line 17

If a corresponding list item of the original list item is in the enclosing device data
environment, the new device data environment uses the corresponding list item from the
enclosing device data environment. No additional storage is allocated in the new device
data environment and neither initialization nor assignment is performed, regardless of
the map-type that is specified.

When running the program below the output I get is:

line 32: Error on "x_start" expected value 2 found value 3
line 37: Error on "x" expected value 6 found value 5

This is with "icc version 14.0.4 (gcc version 4.4.7 compatibility)". The "map(x)" clause is working correctly on the "target" clause in line 27 but not on the "target data" clause in line 24. While it uses the existing "corresponding list item" like it should, it still acts as a "tofrom", updating the corresponding variable before executing the clause's body and updating the original variable after executing the clause's body.

Unless if I missed something in the OpenMP specification the "map" clause of the "target data" directive should be ignored in this case, just like the "map" from the "target" directive does.

Test code:

#include <omp.h>
#include <stdio.h>

#define TEST(var, expected) if (var != expected) { printf("line %d: Error on \"%s\" expected value %d found value %d\n", __LINE__, #var, expected, var); successful = 0; }
char successful = 1;

int main(int argc, char *argv[]) {
    int x = 0, x_start = 0;

    x = 0;
    #pragma omp target data device(1) map(x)
    { //--> x
        x = 1;
        #pragma omp target device(1) map(x)
        { //map should be ignored
            x_start = x;
            x = 2;
        }
        TEST(x_start, 0);
        TEST(x, 1);

        x = 3;

        #pragma omp target data device(1) map(x)
        { //map should be ignored
            x = 4;
            #pragma omp target device(1) map(x)
            { //map should be ignored
                x_start = x;
                x = 5;
            } //map should be ignored
            TEST(x_start, 2);
            TEST(x, 4);

            x = 6;
        } //map should be ignored
        TEST(x, 6);

        x = 7;
    } //<-- x
    TEST(x, 5);

    if(successful)
        printf("Successfully passed all tests\n");
}

[Andrey_C_Intel1]

Alexandros,

You are asking on the wrong forum, because the OpenMP target* constructs are implemented in compiler and are not related to the OpenMP runtime which we are discussing on this forum. Please re-submit your question on the compiler forum so that experts on the OpenMP target construct could answer your question.

Regards,
Andrey

[Alexandros_P_]

My bad, I expected at least the host part of the "target" construct to be in the runtime, I guess I should have checked the source code first. Thanks for me pointing to the correct  forum, I moved the post there.

[Feilong_H_Intel]

Hi Alex,

It is uncommon to have nested #pragma omp target.  And OMP 4.0 Spec says (page 80 line 13):

"If a target, target update, or target data construct appears within a target region then the behavior is unspecified."

I don't have an idea what you want to do exactly in the nested target constructs.  Is it just for testing?

Thanks.

[Alexandros_P_]

I agree that it is uncommon to have nested "target data" constructs and yes the code is just for testing. The line you quoted from the specification applies to "#pragma omp target" not to "#pragma omp target data" and since the specification doesn't state otherwise it is valid code.

You could have nested "target data" clauses targeting different devices, which is why we started testing icc, since it is the only release compiler with a working "target" implimentation. This can also be seen in this discussion http://openmp.org/forum/viewtopic.php?f=13&t=1737

The problem started when we tried to implement the "if" clause in our source to source compiler (which by the way is a bit bugged in icc exactly because icc does not take into account the nested device data environments). The specification says (in both "target" and "target data") that if the "if" clause expression is false, the device is the host and that's why we wanted to see how other compilers impliment that, since the creation of the device data environment has to be dynamic and to happen during runtime as opposed to what we had implimented.

[Ravi_N_Intel]

Yes you can have nested target data clauses.  Only the variables in the outer most target data clauses are transferred.

The current implementation in the compiler has a bug for stack local variables.  They are currently treated map(to/from).  Try using global variables with #omp pragma declare target/#omp pragma end declare target.  But then you would need to use #pragma target update to send and receive data since making it global would mean they are always present.

Thanks