[PATCH] D80051: [OpenMPOpt] Test case 1 - Latency Hiding for Host to Device Memory Transfers

[Hamilton Tobon-Mosquera via Phabricator via llvm-commits]

hamax97 updated this revision to Diff 265929.
hamax97 added a comment.

Added `__tgt_target_data_begin` to the commented original code, just to make the test cases easier to read.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D80051/new/

https://reviews.llvm.org/D80051

Files:
  llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll


Index: llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
===================================================================
--- llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
+++ llvm/test/Transforms/OpenMP/mem_transfer_hiding.ll
@@ -32,7 +32,11 @@
 ;  double random = rand();
 ;  double a = rand() % 777;
 ;
-;  #pragma omp target data map(a)
+;  //#pragma omp target data map(a)
+;  void* args[1];
+;  args[0] = &a;
+;  __tgt_target_data_begin(..., args, ...)
+;
 ;  #pragma omp target teams
 ;  for (int i = 0; i < 1000; ++i) {
 ;    a *= i*i / 2;
@@ -145,7 +149,12 @@
 ;int heavyComputation2(double* a, unsigned size) {
 ;  int random = rand() % 7;
 ;
-;  #pragma omp target data map(a[0:size], size)
+;  //#pragma omp target data map(a[0:size], size)
+;  void* args[2];
+;  args[0] = &a;
+;  args[1] = &size;
+;  __tgt_target_data_begin(..., args, ...)
+;
 ;  #pragma omp target teams
 ;  for (int i = 0; i < size; ++i) {
 ;    a[i] = ++a[i] * 3.141624;
@@ -261,7 +270,12 @@
 ;int heavyComputation3(double* restrict a, unsigned size) {
 ;  int random = rand() % 7;
 ;
-;  #pragma omp target data map(a[0:size], size)
+;  //#pragma omp target data map(a[0:size], size)
+;  void* args[2];
+;  args[0] = &a;
+;  args[1] = &size;
+;  __tgt_target_data_begin(..., args, ...)
+;
 ;  #pragma omp target teams
 ;  for (int i = 0; i < size; ++i) {
 ;    a[i] = ++a[i] * 3.141624;
@@ -414,7 +428,10 @@
 ;  // Random computation.
 ;  int random = rand();
 ;
-;  #pragma omp target data map(to:a[0:size])
+;  //#pragma omp target data map(to:a[0:size])
+;  void* args[1];
+;  args[0] = &a;
+;  __tgt_target_data_begin(..., args, ...)
 ;
 ;  // Random computation.
 ;  random %= size;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D80051.265929.patch
Type: text/x-patch
Size: 1683 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20200524/7eadc249/attachment-0001.bin>