本文介绍了OpenMP中的四个子句(is_device_ptr,use_device_ptr,has_device_addr,use_device_addr)，它们用于在target指令、targetdata和dispatch结构中处理设备内存引用。子句说明了如何指定变量为设备指针，以及如何映射和访问设备上的存储。这些子句在C/C++和Fortran中使用，确保正确管理主机和设备间的内存交互。
本章节翻译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. |
扫一扫
164
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
