oneAPI GPU 优化指南 - 子句: is_device_ptr, use_device_ptr, has_device_addr, use_device

本章节翻译by chenchensmail@163.com 原文：Clauses: is_device_ptr, use_device_ptr, has_device_addr,... (intel.com)

OpenMP 子句 is_device_ptr, use_device_ptr, has_device_addr 和 use_device_addr 可以用来传达在 target、 target data 或 dispatch 结构中引用的变量的信息。这些子句的描述如下。

`is_device_ptr`

is_device_ptr 子句出现在 target 或 dispatch 指令上。它表示列表项是设备指针。因此，在结构体内部对每个列表项进行私有化，并将新的列表项初始化为原始列表项所引用的设备地址。

在 C 语言中，每个列表项应为指针类型或数组类型。
在 C++ 中，每个列表项应为指针类型、数组类型、指针的引用类型或数组的引用类型。
在 Fortran 中，每个列表项应为 C_PTR 类型。

以下 C/C++ 示例说明了如何使用 is_device_ptr 子句。 omp_target_alloc_device 例程在设备上分配内存，并返回该内存的设备指针，该指针保存在主机变量 arr_device 中。在 target 指令上，我们使用 is_device_ptr(arr_device) 子句来表示 arr_device 指向设备内存。因此，在 target 结构体内部， arr_device 被私有化并初始化为 arr_device 所引用的设备地址。

 1#include <stdio.h>
 2#include <stdlib.h>
 3#include <stdint.h>
 4#include <math.h>
 5#include <omp.h>
 6
 7#define N 100
 8
 9int main(void)
10{
11  int *arr_host = NULL;
12  int *arr_device = NULL;
13
14  arr_host = (int *) malloc(N * sizeof(int));
15  arr_device = (int *) omp_target_alloc_device(N * sizeof(int),
16					       omp_get_default_device());
17
18  #pragma omp target is_device_ptr(arr_device) map(from: arr_host[0:N])
19  {
20    for (int i = 0; i < N; ++i) {
21      arr_device[i] = i;
22      arr_host[i] = arr_device[i];
23    }
24  }
25
26  printf ("%d, %d, %d \n", arr_host[0], arr_host[N/2], arr_host[N-1]);
27}

`use_device_ptr`

use_device_ptr 子句出现在一个 target data 指令上。它表示每个列表项都是一个指向在设备上有对应存储或可访问的对象的指针。

如果一个列表项是一个指向映射到设备的对象的指针，那么在结构中对列表项的引用将被转换为对一个设备指针的引用，该设备指针是局部于结构体并且引用相应对象的设备地址。

如果列表项不指向映射对象，它必须包含一个有效的设备地址，列表项引用将被转换为对一个局部设备指针的引用，该指针引用这个设备地址。

每个列表项必须是一个指针，其值是在目标数据环境中具有对应存储或在目标设备上可访问的对象的地址。

在 C 语言中，每个列表项应为指针类型或数组类型。

在 C++ 中，每个列表项应为指针类型、数组类型、指针的引用类型或数组的引用类型。

在 Fortran 中，每个列表项应为 C_PTR 类型。

以下 C/C++ 示例说明了如何使用 use_device_ptr 子句。 omp_target_alloc_device 例程被调用三次以在设备上分配内存。分配的内存的地址保存在主机上的指针变量 A, B 和 C 中。我们使用 use_device_ptr(A, B,C) 子句在 target data 指令上来表示 A, B 和 C 包含有效的设备地址。

 1#include <stdio.h>
 2#include <stdlib.h>
 3#include <stdint.h>
 4#include <math.h>
 5#include <omp.h>
 6
 7#define length 65536
 8
 9int main(void)
10{
11  int device_id = omp_get_default_device();
12  size_t bytes = length*sizeof(double);
13  double * __restrict A;
14  double * __restrict B;
15  double * __restrict C;
16  double scalar = 3.0;
17  double ar;
18  double br;
19  double cr;
20  double asum;
21
22  // Allocate arrays in device memory
23
24  A = (double *) omp_target_alloc_device(bytes, device_id);
25  if (A == NULL){
26     printf(" ERROR: Cannot allocate space for A using omp_target_alloc_device().\n");
27     exit(1);
28  }
29
30  B = (double *) omp_target_alloc_device(bytes, device_id);
31  if (B == NULL){
32     printf(" ERROR: Cannot allocate space for B using omp_target_alloc_device().\n");
33     exit(1);
34  }
35
36  C = (double *) omp_target_alloc_device(bytes, device_id);
37  if (C == NULL){
38     printf(" ERROR: Cannot allocate space for C using omp_target_alloc_device().\n");
39     exit(1);
40  }
41
42  #pragma omp target data use_device_ptr(A,B,C)
43  {
44      // Initialize the arrays
45
46      #pragma omp target teams distribute parallel for
47      for (size_t i=0; i<length; i++) {
48          A[i] = 2.0;
49          B[i] = 2.0;
50          C[i] = 0.0;
51      }
52
53      // Perform the computation
54
55      #pragma omp target teams distribute parallel for
56      for (size_t i=0; i<length; i++) {
57          C[i] += A[i] + scalar * B[i];
58      }
59
60      // Validate and output results
61
62      ar = 2.0;
63      br = 2.0;
64      cr = 0.0;
65      for (int i=0; i<length; i++) {
66          cr += ar + scalar * br;
67      }
68
69      asum = 0.0;
70      #pragma omp target teams distribute parallel for reduction(+:asum)
71      for (size_t i=0; i<length; i++) {
72          asum += fabs(C[i]);
73      }
74
75  } // end target data
76
77  omp_target_free(A, device_id);
78  omp_target_free(B, device_id);
79  omp_target_free(C, device_id);
80
81  double epsilon=1.e-8;
82  if (fabs(cr - asum)/asum > epsilon) {
83      printf("Failed Validation on output array\n"
84             "       Expected checksum: %lf\n"
85             "       Observed checksum: %lf\n"
86             "ERROR: solution did not validate\n", cr, asum);
87      return 1;
88  } else {
89      printf("Solution validates. Checksum = %lf\n", asum);
90  }
91
92  return 0;
93}

`has_device_addr`

has_device_addr 子句出现在一个 target 指令上。它表示列表项已经具有有效的设备地址，因此可以直接从设备访问。

每个列表项必须对于设备数据环境具有有效的设备地址。它可以是任何类型，包括数组部分。

特别地，在 Fortran 中， has_device_addr 子句非常有用，因为它可以用于任何类型的列表项（不仅仅是 C_PTR ）来表示列表项具有设备地址。

以下的 Fortran 示例说明了如何使用 has_device_addr 子句。在示例中，三个数组 A, B 和 C 在设备上分配。当在 target 区域引用数组时，我们使用 has_device_addr(A, B, C) 子句来表示 A, B 和 C 已经具有设备地址。

 1program main
 2  use iso_fortran_env
 3  use omp_lib
 4  implicit none
 5
 6  integer, parameter :: iterations=1000
 7  integer, parameter :: length=64*1024*1024
 8  real(kind=REAL64), parameter ::  epsilon=1.D-8
 9  real(kind=REAL64), allocatable ::  A(:)
10  real(kind=REAL64), allocatable ::  B(:)
11  real(kind=REAL64), allocatable ::  C(:)
12  real(kind=REAL64) :: scalar=3.0
13  real(kind=REAL64) :: ar, br, cr, asum
14  real(kind=REAL64) :: nstream_time, avgtime
15  integer :: i, iter
16
17  !
18  ! Allocate arrays in device memory
19
20  !$omp allocate allocator(omp_target_device_mem_alloc)
21  allocate(A(length))
22
23  !$omp allocate allocator(omp_target_device_mem_alloc)
24  allocate(B(length))
25
26  !$omp allocate allocator(omp_target_device_mem_alloc)
27  allocate(C(length))
28
29  !
30  ! Initialize the arrays
31
32  !$omp target teams distribute parallel do has_device_addr(A, B, C)
33  do i = 1, length
34      A(i) = 2.0
35      B(i) = 2.0
36      C(i) = 0.0
37  end do
38
39  !
40  ! Perform the computation
41
42  nstream_time = omp_get_wtime()
43  do iter = 1, iterations
44     !$omp target teams distribute parallel do has_device_addr(A, B, C)
45     do i = 1, length
46        C(i) = C(i) + A(i) + scalar * B(i)
47     end do
48  end do
49  nstream_time = omp_get_wtime() - nstream_time
50
51  !
52  ! Validate and output results
53
54  ar = 2.0
55  br = 2.0
56  cr = 0.0
57  do iter = 1, iterations
58     do i = 1, length
59        cr = cr + ar + scalar * br
60     end do
61  end do
62
63  asum = 0.0
64  !$omp target teams distribute parallel do reduction(+:asum) has_device_addr(C)
65  do i = 1, length
66     asum = asum + abs(C(i))
67  end do
68
69  if (abs(cr - asum)/asum > epsilon) then
70     print *, "Failed Validation on output array:", "Expected =", cr, "Observed =", asum
71  else
72     avgtime = nstream_time/iterations
73     print *, "Solution validates:", "Checksum =", asum, "Avg time (s) =",  avgtime
74  endif
75
76  deallocate(A)
77  deallocate(B)
78  deallocate(C)
79
80end program main

`use_device_addr`

use_device_addr 子句出现在一个 target data 指令上。它表示每个列表项已经在设备上有对应的存储或者可以在设备上访问。

如果一个列表项被映射，那么在结构中对列表项的引用将被转换为对相应列表项的引用。如果一个列表项没有被映射，那么假定它可以在设备上访问。

一个列表项可以是一个数组部分。

就像 has_device_addr, use_device_addr 子句在 Fortran 中特别有用，因为它可以用于任何类型的列表项（不仅仅是 C_PTR ）来表示列表项具有设备地址。

以下的 Fortran 示例说明了如何使用 use_device_addr 子句。在示例中， array_d 使用 alloc 映射类型映射到设备，因此为 array_d 在设备上分配存储空间，并且不会发生主机和设备之间的数据传输。我们在 target data 指令上使用 use_device_addr(array_d) 子句来表示 array_d 在设备上有对应的存储。

 1program target_use_device_addr
 2
 3  use omp_lib
 4  use iso_fortran_env, only : real64
 5  implicit none
 6
 7  integer, parameter :: N1 = 1024
 8  real(kind=real64), parameter :: aval = real(42, real64)
 9  real(kind=real64), allocatable :: array_d(:), array_h(:)
10  integer :: i,err
11
12  ! Allocate host data
13  allocate(array_h(N1))
14
15  !$omp target data map (from:array_h(1:N1)) map(alloc:array_d(1:N1))
16  !$omp target data use_device_addr(array_d)
17  !$omp target
18      do i=1, N1
19        array_d(i) = aval
20        array_h(i) = array_d(i)
21     end do
22  !$omp end target
23  !$omp end target data
24  !$omp end target data
25
26  ! Check result
27  write (*,*) array_h(1), array_h(N1)
28  if (any(array_h /= aval)) then
29    err = 1
30  else
31    err = 0
32  end if
33
34  deallocate(array_h)
35  if (err == 1) then
36    stop 1
37  else
38    stop 0
39  end if
40
41end program target_use_device_addr

下表总结了本节描述的子句的属性。

Clause	On which directive	Type of list item	Description
is_device_ptr	target, dispatch	C/C++: Pointer, array, or reference Fortran: C_PTR	Indicates that list item is a device pointer (has valid device address).
use_device_ptr	target data	C/C++: Pointer, array, or reference Fortran: C_PTR	Indicates that list item is a pointer to an object that has corresponding storage on device or is accessible on device.
has_device_addr	target	Any type (may be array section)	Indicates that list item has a valid device address.
use_device_addr	target data	Any type (may be array section)	Indicates that list item has corresponding storage on device or is accessible on the device.