Page MenuHomePhabricator

[OpenMP5.0] map item can be non-contiguous for target update
AcceptedPublic

Authored by cchen on Jul 20 2020, 11:06 AM.

Details

Reviewers
ABataev
jdoerfert
Summary

In order not to modify the tgt_target_data_update information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload arg when
the maptype is set as OMP_MAP_DESCRIPTOR. The origin arg is for
passing the pointer information, however, the overloaded arg is an
array of descriptor_dim:

struct descriptor_dim {

int64_t offset;
int64_t count;
int64_t stride

};

and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
arg_size parameter by using dimension size.

For supporting stride in array section, we use a dummy dimension in
descriptor to store the unit size. The formula for counting the stride
in dimension D_n: unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride.

Demonstrate how it works:

double arr[3][4][5];

D0: { offset = 0, count = 1, stride = 8 }                                // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 }                   // stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40  }            // stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 }        // stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200

// X here means we need to offload this data, therefore, runtime will transfer
// data from offset 80, 96, 120, 136, 400, 416, 440, 456
// Runtime patch: https://reviews.llvm.org/D82245
// OOOOO OOOOO OOOOO
// OOOOO OOOOO OOOOO
// XOXOO OOOOO XOXOO
// XOXOO OOOOO XOXOO

Diff Detail

Unit TestsFailed

TimeTest
470 mslinux > MemorySanitizer-X86_64.MemorySanitizer-X86_64::Unknown Unit Message ("")
Script: -- : 'RUN: at line 1'; /mnt/disks/ssd0/agent/llvm-project/build/./bin/clang --driver-mode=g++ -fsanitize=memory -mno-omit-leaf-frame-pointer -fno-omit-frame-pointer -fno-optimize-sibling-calls -m64 -gline-tables-only -std=c++11 -O0 -g /mnt/disks/ssd0/agent/llvm-project/compiler-rt/test/msan/__strxfrm_l.cpp -o /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/X86_64/Output/__strxfrm_l.cpp.tmp && /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/X86_64/Output/__strxfrm_l.cpp.tmp
490 mslinux > MemorySanitizer-X86_64.MemorySanitizer-X86_64::Unknown Unit Message ("")
Script: -- : 'RUN: at line 1'; /mnt/disks/ssd0/agent/llvm-project/build/./bin/clang --driver-mode=g++ -fsanitize=memory -mno-omit-leaf-frame-pointer -fno-omit-frame-pointer -fno-optimize-sibling-calls -m64 -gline-tables-only -O0 -g /mnt/disks/ssd0/agent/llvm-project/compiler-rt/test/msan/strxfrm.cpp -o /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/X86_64/Output/strxfrm.cpp.tmp && /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/X86_64/Output/strxfrm.cpp.tmp
230 mslinux > MemorySanitizer-lld-X86_64.MemorySanitizer-lld-X86_64::Unknown Unit Message ("")
Script: -- : 'RUN: at line 1'; /mnt/disks/ssd0/agent/llvm-project/build/./bin/clang --driver-mode=g++ -fsanitize=memory -mno-omit-leaf-frame-pointer -fno-omit-frame-pointer -fno-optimize-sibling-calls -m64 -fuse-ld=lld -gline-tables-only -std=c++11 -O0 -g /mnt/disks/ssd0/agent/llvm-project/compiler-rt/test/msan/__strxfrm_l.cpp -o /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/lld-X86_64/Output/__strxfrm_l.cpp.tmp && /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/lld-X86_64/Output/__strxfrm_l.cpp.tmp
270 mslinux > MemorySanitizer-lld-X86_64.MemorySanitizer-lld-X86_64::Unknown Unit Message ("")
Script: -- : 'RUN: at line 1'; /mnt/disks/ssd0/agent/llvm-project/build/./bin/clang --driver-mode=g++ -fsanitize=memory -mno-omit-leaf-frame-pointer -fno-omit-frame-pointer -fno-optimize-sibling-calls -m64 -fuse-ld=lld -gline-tables-only -O0 -g /mnt/disks/ssd0/agent/llvm-project/compiler-rt/test/msan/strxfrm.cpp -o /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/lld-X86_64/Output/strxfrm.cpp.tmp && /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/msan/lld-X86_64/Output/strxfrm.cpp.tmp
520 mslinux > SanitizerCommon-asan-x86_64-Linux.Linux::Unknown Unit Message ("")
Script: -- : 'RUN: at line 1'; /mnt/disks/ssd0/agent/llvm-project/build/./bin/clang --driver-mode=g++ -gline-tables-only -fsanitize=address -m64 -ldl -std=c++11 -O0 -g /mnt/disks/ssd0/agent/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/protoent.cpp -o /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/protoent.cpp.tmp && /mnt/disks/ssd0/agent/llvm-project/build/projects/compiler-rt/test/sanitizer_common/asan-x86_64-Linux/Linux/Output/protoent.cpp.tmp 2>&1 | FileCheck /mnt/disks/ssd0/agent/llvm-project/compiler-rt/test/sanitizer_common/TestCases/Linux/protoent.cpp
View Full Test Results (14 Failed)

Event Timeline

cchen created this revision.Jul 20 2020, 11:06 AM

Also, what about compatibility with declare mapper? Can you add tests for it?

clang/lib/CodeGen/CGOpenMPRuntime.cpp
7809–7813

DO not use braced initializer list https://llvm.org/docs/CodingStandards.html#id26

7839–7847

No need for braces here per coding standard

7876

have already have

7892

debug code

7893

Move this after the first if statement

7894–7896

Better to make it something like:

if (const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr))
7944–7945

If no stride at all, why need to divide?

7967

Same, if no stride at all, no need to mul.

clang/lib/Sema/SemaOpenMP.cpp
16913–16919

Better to transform it to something like:

if (!OASE || OASE->getLength())
  break;
SemaRef.Diag(...)
cchen added a comment.Jul 21 2020, 2:03 PM

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

cchen updated this revision to Diff 279631.Jul 21 2020, 2:04 PM

Fix based on feedback

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list
cchen added a comment.Jul 21 2020, 2:20 PM

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

Oh, I didn't see that. I'll add test for it, thanks!

cchen added a comment.Jul 21 2020, 2:59 PM

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

cchen added a comment.Jul 21 2020, 3:52 PM

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

After discussing with @dreachem, my understanding is that it is not incorrect to not allowing extended/non-contiguous array section to appear in declare mapper.
For the declare mapper directive, since spec only allows map clause, extended array section (with stride) or non-contiguous array section are both not allowed.
For using the mapper in map/to/from clause, if mapping or updating an array section of type T, where there is a mapper declared for T, then the mapper needs to apply as if to each element of the array or array section. Spec is intentionally not sufficiently clear on this for target update so the semantic check in Clang is not incorrect. Which lead to the fact that I might not need to support extended/non-contiguous array section for declare mapper.

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

After discussing with @dreachem, my understanding is that it is not incorrect to not allowing extended/non-contiguous array section to appear in declare mapper.
For the declare mapper directive, since spec only allows map clause, extended array section (with stride) or non-contiguous array section are both not allowed.
For using the mapper in map/to/from clause, if mapping or updating an array section of type T, where there is a mapper declared for T, then the mapper needs to apply as if to each element of the array or array section. Spec is intentionally not sufficiently clear on this for target update so the semantic check in Clang is not incorrect. Which lead to the fact that I might not need to support extended/non-contiguous array section for declare mapper.

What exactly is incorrect in clang sema?

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

After discussing with @dreachem, my understanding is that it is not incorrect to not allowing extended/non-contiguous array section to appear in declare mapper.
For the declare mapper directive, since spec only allows map clause, extended array section (with stride) or non-contiguous array section are both not allowed.
For using the mapper in map/to/from clause, if mapping or updating an array section of type T, where there is a mapper declared for T, then the mapper needs to apply as if to each element of the array or array section. Spec is intentionally not sufficiently clear on this for target update so the semantic check in Clang is not incorrect. Which lead to the fact that I might not need to support extended/non-contiguous array section for declare mapper.

What exactly is incorrect in clang sema?

I mean "not" incorrect for not allowing extended/non-contiguous array section in `to/from([mapper(mapper-identifier):] location-list).

cchen marked 9 inline comments as done.Jul 23 2020, 11:40 AM
cchen added a comment.Jul 24 2020, 8:36 AM

@ABataev, is there any other concern for this patch?

ABataev added inline comments.Jul 24 2020, 8:43 AM
clang/lib/CodeGen/CGOpenMPRuntime.cpp
8914
  1. The second argument must be of integer type, not boolean.
  2. Why it is signed?
clang/lib/Sema/SemaOpenMP.cpp
16915

No need for else here

cchen updated this revision to Diff 280474.Jul 24 2020, 8:54 AM

Fix coding style and argument on getIntTypeForBitwidth

This revision is now accepted and ready to land.Jul 24 2020, 9:02 AM
cchen added a comment.Jul 24 2020, 9:06 AM

@ABataev, do I need to wait for the runtime patch to commit this one? If so, do you have some recommend reviewers for me to add to that patch? I have pinged several times for that patch but haven't got many reviews for it. Thanks!

@ABataev, do I need to wait for the runtime patch to commit this one? If so, do you have some recommend reviewers for me to add to that patch? I have pinged several times for that patch but haven't got many reviews for it. Thanks!

Yes, it must be committed after runtime patch.