LI80943: ISSUES WHEN USING OMP_TARGET_MEMCPY()

APAR status

• Closed as program error.

Error description

• omp_target_memcpy() does not work when both
  source and destination are device pointers (i.e.
  device-to-device copy) as shown in the following testcase:
  
  ===== COMPILE COMMAND:
  xlc -qsmp=omp -qoffload test.c
  
  ===== TESTCASE:
  $ cat test.c
  #include <omp.h>
  #include <stdio.h>
  #include <stdlib.h>
  
  // #include <cuda.h>
  
  int main()
    {
    int nValues = 10;
    double *Value_1, *Value_2;
    void *D_Value_1, *D_Value_2;
  
    Value_1 = malloc ( nValues * sizeof ( double ) );
    Value_2 = malloc ( nValues * sizeof ( double ) );
  
    for ( int iV = 1; iV < nValues; iV++ )
      Value_1 [ iV ] = ( ( double ) rand() ) /  ( ( double )
  RAND_MAX );
  
    #pragma omp target enter data \
       map ( to: Value_1 [ : nValues ] ) \
       map ( alloc:  Value_2 [ : nValues ] )
  
    printf ( "Value_1: %f, Value_2: %f\n", Value_1 [ 5 ], Value_2
  [ 5 ] );
  
    printf ("host addr %p, %p\n", Value_1, Value_2);
    #pragma omp target data map ( nValues ) use_device_ptr (
  Value_1, Value_2 )
    {
    D_Value_1 = Value_1;
    D_Value_2 = Value_2;
    }
  
    printf ("Device addr %p, %p\n", D_Value_1, D_Value_2);
    printf ("Host addr %p, %p\n", Value_1, Value_2);
    printf ("OMP default device: %d\n", omp_get_default_device ( )
  );
    printf ("OMP init device: %d\n", omp_get_initial_device ( ) );
  
    int res
      = omp_target_memcpy
           ( D_Value_2, D_Value_1, nValues * sizeof ( double ), 0,
  0,
             omp_get_default_device ( ), omp_get_default_device (
  ) );
  
    // cudaMemcpy ( D_Value_2, D_Value_1, nValues * sizeof (
  double ), 3 );
  
    // printf ( "res: %d\n", res );
  
    for ( int iV = 1; iV < nValues; iV++ )
      Value_1 [ iV ] = 0.0;
  
    printf ( "Before Value_1: %f, Value_2: %f\n", Value_1 [ 5 ],
  Value_2 [ 5 ]
  );
  
    #pragma omp target update from \
        ( Value_2 [ : nValues ], Value_1 [ : nValues ] )
  
    printf ( "After Value_1: %f, Value_2: %f\n", Value_1 [ 5 ],
  Value_2 [ 5 ]
  );
  
  
    printf("Hello, World!\n");
    return 0;
    }
  
  $
  
  
  ===== ACTUAL OUTPUT:
  $ ./a.out
  Value_1: 0.911647, Value_2: 0.000000
  host addr 0x1ce90330, 0x1ce90390
  Device addr 0x7fff20000000, 0x7fff20000080
  Host addr 0x1ce90330, 0x1ce90390
  OMP default device: 0
  OMP init device: -10
  Before Value_1: 0.000000, Value_2: 0.000000
  After Value_1: 0.911647, Value_2: 0.000000
  Hello, World!
  $
  
  The above is not expected output. After Value_1 does not match
  After Value_2. This is to represent that the omp_target_memcpy()
  call on lines 37-40 is unsuccessful when both source and
  destination are device pointers (D_Value_[1,2]).
  
  
  
  ===== EXPECTED OUTPUT:
  $ ./a.out
  Value_1: 0.911647, Value_2: 0.000000
  host addr 0x4a6c0330, 0x4a6c0390
  Device addr 0x7fff20000000, 0x7fff20000080
  Host addr 0x4a6c0330, 0x4a6c0390
  OMP default device: 0
  OMP init device: -10
  Before Value_1: 0.000000, Value_2: 0.000000
  After Value_1: 0.911647, Value_2: 0.911647
  Hello, World!
  $
  
  The above is the expected output. After Value_1 matches After
  Value_2. This is to represent that the omp_target_memcpy() call
  on lines 37-40 is successful when both source and destination
  are device pointers (D_Value_[1,2]).
  

Local fix

• N/A
  

Problem summary

• USERS AFFECTED:
  Users who use omp_target_memcpy() to copy data from device to
  device may be affected by this issue.
  
  PROBLEM DESCRIPTION:
  When the OMP Runtime retrieves the data on the device, the
  addresses of the source and destination were incorrect.
  

Problem conclusion

• This problem has been fixed.
  

Temporary fix

Comments

APAR Information

• APAR is sysrouted FROM one or more of the following:

• APAR is sysrouted TO one or more of the following:

Fix information

• Fixed component name

  XL C/C++ LINUX

• Fixed component ID

  5725C7310

Applicable component levels