From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id 1BDA3384D1B2; Wed, 29 Jun 2022 14:46:11 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 1BDA3384D1B2 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-12] openmp: Do USM transform for omp_target_alloc X-Act-Checkin: gcc X-Git-Author: Andrew Stubbs X-Git-Refname: refs/heads/devel/omp/gcc-12 X-Git-Oldrev: 8099ba0e264b1292a785417381947f274616b5b0 X-Git-Newrev: 39bdcd849ec435bd8349161e7ad3e2c94160f645 Message-Id: <20220629144611.1BDA3384D1B2@sourceware.org> Date: Wed, 29 Jun 2022 14:46:11 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Wed, 29 Jun 2022 14:46:11 -0000 https://gcc.gnu.org/g:39bdcd849ec435bd8349161e7ad3e2c94160f645 commit 39bdcd849ec435bd8349161e7ad3e2c94160f645 Author: Andrew Stubbs Date: Fri Apr 1 16:12:16 2022 +0100 openmp: Do USM transform for omp_target_alloc OpenMP 5.0 says that omp_target_alloc should return USM addresses. gcc/ChangeLog: * omp-low.cc (usm_transform): Transform omp_target_alloc and omp_target_free. libgomp/ChangeLog: * testsuite/libgomp.c/usm-6.c: Add omp_target_alloc. gcc/testsuite/ChangeLog: * c-c++-common/gomp/usm-2.c: Add omp_target_alloc. * c-c++-common/gomp/usm-3.c: Add omp_target_alloc. Diff: --- gcc/ChangeLog.omp | 5 +++++ gcc/omp-low.cc | 6 ++++-- gcc/testsuite/ChangeLog.omp | 5 +++++ gcc/testsuite/c-c++-common/gomp/usm-2.c | 9 ++++++++- gcc/testsuite/c-c++-common/gomp/usm-3.c | 9 ++++++++- libgomp/ChangeLog.omp | 4 ++++ libgomp/testsuite/libgomp.c/usm-6.c | 13 +++++++++++-- 7 files changed, 45 insertions(+), 6 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index c419396e16f..a98d8b3cfed 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,8 @@ +2022-04-02 Andrew Stubbs + + * omp-low.cc (usm_transform): Transform omp_target_alloc and + omp_target_free. + 2022-03-31 Abid Qadeer * omp-low.cc (lower_omp_allocate): Move allocate declaration diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 8685a0b6abc..0734a4e9743 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -16143,7 +16143,8 @@ usm_transform (gimple_stmt_iterator *gsi_p, bool *, if ((strcmp (name, "malloc") == 0) || (fndecl_built_in_p (fndecl, BUILT_IN_NORMAL) && DECL_FUNCTION_CODE (fndecl) == BUILT_IN_MALLOC) - || DECL_IS_REPLACEABLE_OPERATOR_NEW_P (fndecl)) + || DECL_IS_REPLACEABLE_OPERATOR_NEW_P (fndecl) + || strcmp (name, "omp_target_alloc") == 0) { tree omp_alloc_type = build_function_type_list (ptr_type_node, size_type_node, @@ -16215,7 +16216,8 @@ usm_transform (gimple_stmt_iterator *gsi_p, bool *, || (fndecl_built_in_p (fndecl, BUILT_IN_NORMAL) && DECL_FUNCTION_CODE (fndecl) == BUILT_IN_FREE) || (DECL_IS_OPERATOR_DELETE_P (fndecl) - && DECL_IS_REPLACEABLE_OPERATOR (fndecl))) + && DECL_IS_REPLACEABLE_OPERATOR (fndecl)) + || strcmp (name, "omp_target_free") == 0) { tree omp_free_type = build_function_type_list (void_type_node, ptr_type_node, diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 6ff2e7339a8..9d8026e0006 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,8 @@ +2022-04-02 Andrew Stubbs + + * c-c++-common/gomp/usm-2.c: Add omp_target_alloc. + * c-c++-common/gomp/usm-3.c: Add omp_target_alloc. + 2022-03-11 Andrew Stubbs Backport of the patch posted at diff --git a/gcc/testsuite/c-c++-common/gomp/usm-2.c b/gcc/testsuite/c-c++-common/gomp/usm-2.c index 64dbb6be131..8c20ef94e69 100644 --- a/gcc/testsuite/c-c++-common/gomp/usm-2.c +++ b/gcc/testsuite/c-c++-common/gomp/usm-2.c @@ -12,6 +12,8 @@ void *aligned_alloc (__SIZE_TYPE__, __SIZE_TYPE__); void *calloc(__SIZE_TYPE__, __SIZE_TYPE__); void *realloc(void *, __SIZE_TYPE__); void free (void *); +void *omp_target_alloc (__SIZE_TYPE__, int); +void omp_target_free (void *, int); #ifdef __cplusplus } @@ -24,16 +26,21 @@ foo () void *p2 = realloc(p1, 30); void *p3 = calloc(4, 15); void *p4 = aligned_alloc(16, 40); + void *p5 = omp_target_alloc(50, 1); free (p2); free (p3); free (p4); + omp_target_free (p5, 1); } /* { dg-final { scan-tree-dump-times "omp_alloc \\(20, 10\\)" 1 "usm_transform" } } */ /* { dg-final { scan-tree-dump-times "omp_realloc \\(.*, 30, 10, 10\\)" 1 "usm_transform" } } */ /* { dg-final { scan-tree-dump-times "omp_calloc \\(4, 15, 10\\)" 1 "usm_transform" } } */ /* { dg-final { scan-tree-dump-times "omp_aligned_alloc \\(16, 40, 10\\)" 1 "usm_transform" } } */ -/* { dg-final { scan-tree-dump-times "omp_free" 3 "usm_transform" } } */ +/* { dg-final { scan-tree-dump-times "omp_alloc \\(50, 10\\)" 1 "usm_transform" } } */ +/* { dg-final { scan-tree-dump-times "omp_free" 4 "usm_transform" } } */ /* { dg-final { scan-tree-dump-not " free" "usm_transform" } } */ /* { dg-final { scan-tree-dump-not " aligned_alloc" "usm_transform" } } */ /* { dg-final { scan-tree-dump-not " malloc" "usm_transform" } } */ +/* { dg-final { scan-tree-dump-not " omp_target_alloc" "usm_transform" } } */ +/* { dg-final { scan-tree-dump-not " omp_target_free" "usm_transform" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/usm-3.c b/gcc/testsuite/c-c++-common/gomp/usm-3.c index 934582ea5fd..2b0cbb45e27 100644 --- a/gcc/testsuite/c-c++-common/gomp/usm-3.c +++ b/gcc/testsuite/c-c++-common/gomp/usm-3.c @@ -10,6 +10,8 @@ void *aligned_alloc (__SIZE_TYPE__, __SIZE_TYPE__); void *calloc(__SIZE_TYPE__, __SIZE_TYPE__); void *realloc(void *, __SIZE_TYPE__); void free (void *); +void *omp_target_alloc (__SIZE_TYPE__, int); +void omp_target_free (void *, int); #ifdef __cplusplus } @@ -22,16 +24,21 @@ foo () void *p2 = realloc(p1, 30); void *p3 = calloc(4, 15); void *p4 = aligned_alloc(16, 40); + void *p5 = omp_target_alloc(50, 1); free (p2); free (p3); free (p4); + omp_target_free (p5, 1); } /* { dg-final { scan-tree-dump-times "omp_alloc \\(20, 10\\)" 1 "usm_transform" } } */ /* { dg-final { scan-tree-dump-times "omp_realloc \\(.*, 30, 10, 10\\)" 1 "usm_transform" } } */ /* { dg-final { scan-tree-dump-times "omp_calloc \\(4, 15, 10\\)" 1 "usm_transform" } } */ /* { dg-final { scan-tree-dump-times "omp_aligned_alloc \\(16, 40, 10\\)" 1 "usm_transform" } } */ -/* { dg-final { scan-tree-dump-times "omp_free" 3 "usm_transform" } } */ +/* { dg-final { scan-tree-dump-times "omp_alloc \\(50, 10\\)" 1 "usm_transform" } } */ +/* { dg-final { scan-tree-dump-times "omp_free" 4 "usm_transform" } } */ /* { dg-final { scan-tree-dump-not " free" "usm_transform" } } */ /* { dg-final { scan-tree-dump-not " aligned_alloc" "usm_transform" } } */ /* { dg-final { scan-tree-dump-not " malloc" "usm_transform" } } */ +/* { dg-final { scan-tree-dump-not " omp_target_alloc" "usm_transform" } } */ +/* { dg-final { scan-tree-dump-not " omp_target_free" "usm_transform" } } */ diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 888fc16bc03..1693c212a12 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,7 @@ +2022-04-02 Andrew Stubbs + + * testsuite/libgomp.c/usm-6.c: Add omp_target_alloc. + 2022-03-31 Abid Qadeer * testsuite/libgomp.fortran/allocate-2.f90: Remove commented lines. diff --git a/libgomp/testsuite/libgomp.c/usm-6.c b/libgomp/testsuite/libgomp.c/usm-6.c index d2c828fdc9d..c207140092a 100644 --- a/libgomp/testsuite/libgomp.c/usm-6.c +++ b/libgomp/testsuite/libgomp.c/usm-6.c @@ -4,6 +4,8 @@ #include #include +#include + /* On old systems, the declaraition may not be present in stdlib.h which will generate a warning. This function is going to be replaced with omp_aligned_alloc so the purpose of this declaration is to avoid that @@ -19,7 +21,8 @@ main () int *b = (int *) calloc(sizeof(int), 3); int *c = (int *) realloc(NULL, sizeof(int) * 4); int *d = (int *) aligned_alloc(32, sizeof(int)); - if (!a || !b || !c || !d) + int *e = (int *) omp_target_alloc(sizeof(int), 1); + if (!a || !b || !c || !d || !e) __builtin_abort (); a[0] = 42; @@ -36,6 +39,7 @@ main () uintptr_t b_p = (uintptr_t)b; uintptr_t c_p = (uintptr_t)c; uintptr_t d_p = (uintptr_t)d; + uintptr_t e_p = (uintptr_t)e; if (d_p & 31 != 0) __builtin_abort (); @@ -52,9 +56,12 @@ main () __builtin_abort (); if (d_p != (uintptr_t)d) __builtin_abort (); + if (e_p != (uintptr_t)e) + __builtin_abort (); a[0] = 72; b[0] = 82; c[0] = 92; + e[0] = 102; } #pragma omp target @@ -74,10 +81,12 @@ main () if (a[0] != 72 || a[1] != 73 || b[0] != 82 || b[1] != 83 - || c[0] != 92 || c[1] != 93) + || c[0] != 92 || c[1] != 93 + || e[0] != 102) __builtin_abort (); free(a); free(b); free(c); + omp_target_free(e, 1); return 0; }