Skip to content

Commit 2074e13

Browse files
amitamd7amitamd7
authored
[Clang][OpenMP] Non-contiguous strided update (#144635)
This patch handles the strided update in the `#pragma omp target update from(data[a:b:c])` directive where 'c' represents the strided access leading to non-contiguous update in the `data` array when the offloaded execution returns the control back to host from device using the `from` clause. Issue: Clang CodeGen where info is generated for the particular `MapType` (to, from, etc), it was failing to detect the strided access. Because of this, the `MapType` bits were incorrect when passed to runtime. This led to incorrect execution (contiguous) in the libomptarget runtime code. Added a minimal testcase that verifies the working of the patch.
1 parent 72b53cd commit 2074e13

File tree

9 files changed

+323
-2
lines changed

9 files changed

+323
-2
lines changed

‎clang/docs/OpenMPSupport.rst‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -191,7 +191,7 @@ implementation.
191191
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
192192
| device | teams construct on the host device | :good:`done` | r371553 |
193193
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
194-
| device | support non-contiguous array sections for target update | :good:`done` | |
194+
| device | support non-contiguous array sections for target update | :good:`done` | https://github.com/llvm/llvm-project/pull/144635 |
195195
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
196196
| device | pointer attachment | :good:`done` | |
197197
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+

‎clang/docs/ReleaseNotes.rst‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -322,6 +322,7 @@ OpenMP Support
322322
- Added parsing and semantic analysis support for the ``need_device_addr``
323323
modifier in the ``adjust_args`` clause.
324324
- Allow array length to be omitted in array section subscript expression.
325+
- Fixed non-contiguous strided update in the ``omp target update`` directive with the ``from`` clause.
325326

326327
Improvements
327328
^^^^^^^^^^^^

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp‎

Lines changed: 26 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7494,7 +7494,32 @@ class MappableExprsHandler {
74947494
// dimension.
74957495
uint64_t DimSize = 1;
74967496

7497-
bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
7497+
// Detects non-contiguous updates due to strided accesses.
7498+
// Sets the 'IsNonContiguous' flag so that the 'MapType' bits are set
7499+
// correctly when generating information to be passed to the runtime. The
7500+
// flag is set to true if any array section has a stride not equal to 1, or
7501+
// if the stride is not a constant expression (conservatively assumed
7502+
// non-contiguous).
7503+
bool IsNonContiguous =
7504+
CombinedInfo.NonContigInfo.IsNonContiguous ||
7505+
any_of(Components, [&](const auto &Component) {
7506+
const auto *OASE =
7507+
dyn_cast<ArraySectionExpr>(Component.getAssociatedExpression());
7508+
if (!OASE)
7509+
return false;
7510+
7511+
const Expr *StrideExpr = OASE->getStride();
7512+
if (!StrideExpr)
7513+
return false;
7514+
7515+
const auto Constant =
7516+
StrideExpr->getIntegerConstantExpr(CGF.getContext());
7517+
if (!Constant)
7518+
return false;
7519+
7520+
return !Constant->isOne();
7521+
});
7522+
74987523
bool IsPrevMemberReference = false;
74997524

75007525
bool IsPartialMapped =

‎clang/test/OpenMP/target_update_strided_messages.c‎

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
2+
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
3+
4+
void foo(void) {}
5+
6+
int main(int argc, char **argv) {
7+
int len = 8;
8+
double data[len];
9+
10+
// Valid strided array sections
11+
#pragma omp target update from(data[0:4:2]) // OK
12+
{}
13+
14+
#pragma omp target update to(data[0:len/2:2]) // OK
15+
{}
16+
17+
#pragma omp target update from(data[1:3:2]) // OK
18+
{}
19+
20+
// Missing stride (default = 1)
21+
#pragma omp target update from(data[0:4]) // OK
22+
{}
23+
24+
// Invalid stride expressions
25+
#pragma omp target update from(data[0:4:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
26+
27+
#pragma omp target update from(data[0:4:-1]) // expected-error {{section stride is evaluated to a non-positive value -1}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
28+
29+
// Missing colon
30+
#pragma omp target update from(data[0:4 2]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
31+
{}
32+
33+
// Too many colons
34+
#pragma omp target update from(data[0:4:2:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
35+
{}
36+
37+
return 0;
38+
}

‎clang/test/OpenMP/target_update_strided_multiple_messages.c‎

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
2+
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
3+
4+
void foo(void) {}
5+
6+
typedef struct {
7+
int len;
8+
double data[12];
9+
} S;
10+
11+
int main(int argc, char **argv) {
12+
int len = 12;
13+
double data1[len], data2[len];
14+
S s;
15+
16+
// Valid multiple strided array sections
17+
#pragma omp target update from(data1[0:4:2], data2[0:2:5]) // OK
18+
{}
19+
20+
#pragma omp target update to(data1[1:2:3], data2[2:3:2]) // OK
21+
{}
22+
23+
// Mixed strided and regular array sections
24+
#pragma omp target update from(data1[0:len], data2[0:4:2]) // OK
25+
{}
26+
27+
// Struct member arrays with strides
28+
#pragma omp target update from(s.data[0:4:2]) // OK
29+
{}
30+
31+
#pragma omp target update from(s.data[0:s.len/2:2]) // OK
32+
{}
33+
34+
// Invalid stride in one of multiple sections
35+
#pragma omp target update from(data1[0:3:4], data2[0:2:0]) // expected-error {{section stride is evaluated to a non-positive value 0}}
36+
37+
// Complex expressions in multiple arrays
38+
int stride1 = 2, stride2 = 3;
39+
#pragma omp target update from(data1[0:len/2:stride1], data2[1:len/3:stride2]) // OK
40+
{}
41+
42+
// Missing colon
43+
#pragma omp target update from(data1[0:4:2], data2[0:3 4]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
44+
45+
return 0;
46+
}

‎clang/test/OpenMP/target_update_strided_partial_messages.c‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wuninitialized
2+
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wuninitialized
3+
4+
void foo(void) {}
5+
6+
int main(int argc, char **argv) {
7+
int len = 11;
8+
double data[len];
9+
10+
// Valid partial strided updates
11+
#pragma omp target update from(data[0:4:3]) // OK
12+
{}
13+
14+
// Stride larger than length
15+
#pragma omp target update from(data[0:2:10]) // OK
16+
{}
17+
18+
// Valid: complex expressions
19+
int offset = 1;
20+
int count = 3;
21+
int stride = 2;
22+
#pragma omp target update from(data[offset:count:stride]) // OK
23+
{}
24+
25+
// Invalid stride expressions
26+
#pragma omp target update from(data[0:4:offset-1]) // OK if offset > 1
27+
{}
28+
29+
#pragma omp target update from(data[0:count:0]) // expected-error {{section stride is evaluated to a non-positive value 0}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
30+
31+
return 0;
32+
}

‎offload/test/offloading/strided_multiple_update.c‎

Lines changed: 62 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,62 @@
1+
// This test checks that #pragma omp target update from(data1[0:3:4],
2+
// data2[0:2:5]) correctly updates disjoint strided sections of multiple arrays
3+
// from the device to the host.
4+
5+
// RUN: %libomptarget-compile-run-and-check-generic
6+
#include <omp.h>
7+
#include <stdio.h>
8+
9+
int main() {
10+
int len = 12;
11+
double data1[len], data2[len];
12+
13+
// Initial values
14+
#pragma omp target map(tofrom : data1[0 : len], data2[0 : len])
15+
{
16+
for (int i = 0; i < len; i++) {
17+
data1[i] = i;
18+
data2[i] = i * 10;
19+
}
20+
}
21+
22+
printf("original host array values:\n");
23+
printf("data1: ");
24+
for (int i = 0; i < len; i++)
25+
printf("%.1f ", data1[i]);
26+
printf("\ndata2: ");
27+
for (int i = 0; i < len; i++)
28+
printf("%.1f ", data2[i]);
29+
printf("\n\n");
30+
31+
#pragma omp target data map(to : data1[0 : len], data2[0 : len])
32+
{
33+
// Modify arrays on device
34+
#pragma omp target
35+
{
36+
for (int i = 0; i < len; i++)
37+
data1[i] += i;
38+
for (int i = 0; i < len; i++)
39+
data2[i] += 100;
40+
}
41+
42+
// data1[0:3:4] // indices 0,4,8
43+
// data2[0:2:5] // indices 0,5
44+
#pragma omp target update from(data1[0 : 3 : 4], data2[0 : 2 : 5])
45+
}
46+
47+
printf("device array values after update from:\n");
48+
printf("data1: ");
49+
for (int i = 0; i < len; i++)
50+
printf("%.1f ", data1[i]);
51+
printf("\ndata2: ");
52+
for (int i = 0; i < len; i++)
53+
printf("%.1f ", data2[i]);
54+
printf("\n\n");
55+
56+
// CHECK: data1: 0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0 8.0 9.0 10.0 11.0
57+
// CHECK: data2: 0.0 10.0 20.0 30.0 40.0 50.0 60.0 70.0 80.0 90.0 100.0 110.0
58+
59+
// CHECK: data1: 0.0 1.0 2.0 3.0 8.0 5.0 6.0 7.0 16.0 9.0 10.0 11.0
60+
// CHECK: data2: 100.0 10.0 20.0 30.0 40.0 150.0 60.0 70.0 80.0 90.0 100.0
61+
// 110.0
62+
}

‎offload/test/offloading/strided_partial_update.c‎

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
// This test checks that #pragma omp target update from(data[0:4:3]) correctly
2+
// updates every third element (stride 3) from the device to the host, partially
3+
// across the array
4+
5+
// RUN: %libomptarget-compile-run-and-check-generic
6+
#include <omp.h>
7+
#include <stdio.h>
8+
9+
int main() {
10+
int len = 11;
11+
double data[len];
12+
13+
#pragma omp target map(tofrom : data[0 : len])
14+
{
15+
for (int i = 0; i < len; i++)
16+
data[i] = i;
17+
}
18+
19+
// Initial values
20+
printf("original host array values:\n");
21+
for (int i = 0; i < len; i++)
22+
printf("%f\n", data[i]);
23+
printf("\n");
24+
25+
#pragma omp target data map(to : data[0 : len])
26+
{
27+
// Modify arrays on device
28+
#pragma omp target
29+
for (int i = 0; i < len; i++)
30+
data[i] += i;
31+
32+
#pragma omp target update from(data[0 : 4 : 3]) // indices 0,3,6,9
33+
}
34+
35+
printf("device array values after update from:\n");
36+
for (int i = 0; i < len; i++)
37+
printf("%f\n", data[i]);
38+
printf("\n");
39+
40+
// CHECK: 0.000000
41+
// CHECK: 1.000000
42+
// CHECK: 2.000000
43+
// CHECK: 3.000000
44+
// CHECK: 4.000000
45+
// CHECK: 5.000000
46+
// CHECK: 6.000000
47+
// CHECK: 7.000000
48+
// CHECK: 8.000000
49+
// CHECK: 9.000000
50+
// CHECK: 10.000000
51+
52+
// CHECK: 0.000000
53+
// CHECK: 1.000000
54+
// CHECK: 2.000000
55+
// CHECK: 6.000000
56+
// CHECK: 4.000000
57+
// CHECK: 5.000000
58+
// CHECK: 12.000000
59+
// CHECK: 7.000000
60+
// CHECK: 8.000000
61+
// CHECK: 18.000000
62+
// CHECK: 10.000000
63+
}

‎offload/test/offloading/strided_update.c‎

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
// This test checks that "update from" clause in OpenMP is supported when the
2+
// elements are updated in a non-contiguous manner. This test checks that
3+
// #pragma omp target update from(data[0:4:2]) correctly updates only every
4+
// other element (stride 2) from the device to the host
5+
6+
// RUN: %libomptarget-compile-run-and-check-generic
7+
#include <omp.h>
8+
#include <stdio.h>
9+
10+
int main() {
11+
int len = 8;
12+
double data[len];
13+
#pragma omp target map(tofrom : len, data[0 : len])
14+
{
15+
for (int i = 0; i < len; i++) {
16+
data[i] = i;
17+
}
18+
}
19+
// Initial values
20+
printf("original host array values:\n");
21+
for (int i = 0; i < len; i++)
22+
printf("%f\n", data[i]);
23+
printf("\n");
24+
25+
#pragma omp target data map(to : len, data[0 : len])
26+
{
27+
// Modify arrays on device
28+
#pragma omp target
29+
for (int i = 0; i < len; i++) {
30+
data[i] += i;
31+
}
32+
33+
#pragma omp target update from(data[0 : 4 : 2])
34+
}
35+
// CHECK: 0.000000
36+
// CHECK: 1.000000
37+
// CHECK: 4.000000
38+
// CHECK: 3.000000
39+
// CHECK: 8.000000
40+
// CHECK: 5.000000
41+
// CHECK: 12.000000
42+
// CHECK: 7.000000
43+
// CHECK-NOT: 2.000000
44+
// CHECK-NOT: 6.000000
45+
// CHECK-NOT: 10.000000
46+
// CHECK-NOT: 14.000000
47+
48+
printf("from target array results:\n");
49+
for (int i = 0; i < len; i++)
50+
printf("%f\n", data[i]);
51+
printf("\n");
52+
53+
return 0;
54+
}

0 commit comments

Comments
 (0)