本文探讨了 OpenMP Runtime 内存分配方式，包括主机分配、设备分配、共享分配和共享系统分配，介绍了 Intel 编译器支持的内存分配例程及使用子句的示例，还分析了不同分配方式的性能，指出设备分配性能最佳，主机分配性能较差。
本章节翻译by chenchensmail@163.com 原文：
本节将探讨各种分配内存的方式，以及支持的分配类型。主机上的指针与设备上的指针大小相同。
主机分配 由主机拥有，预计将从系统内存中分配。主机分配可以由主机和所有支持的设备访问。 因此，可以在主机和所有支持的设备上使用指向主机分配的相同指针。 预计主机分配不会在系统内存和设备本地内存之间迁移。当在设备上访问指向主机分配的指针时， 数据通常会通过连接设备和主机的总线（如 PCI-Express ）发送。
设备分配 由特定设备拥有，预计将从设备本地内存中分配。 在该设备上可以读取和写入分配的存储空间，但不能直接从主机或任何其他支持的设备访问。
共享分配 可以由主机和所有支持的设备访问。因此，可以在主机和所有支持的设备上使用指向共享分配的相同指针， 就像在主机分配中一样。然而，共享分配并不由任何特定设备拥有，而是预计会在主机和一个或多个设备之间迁移。 这意味着在迁移发生后，在设备上的访问会从更快的设备本地内存中进行，而不是通过高延迟的总线连接远程访问系统内存。
共享系统分配 是共享分配的子类，其中内存由系统分配器（如 malloc 或 new ）而不是分配 API （如 OpenMP 内存分配 API ）分配。 共享系统分配没有关联的设备；它们本质上是跨设备的。像其他共享分配一样，共享系统分配预计会在主机和支持的设备之间迁移， 并且可以在主机和所有支持的设备上使用指向共享系统分配的相同指针。
注意事项:
-
目前，在 Intel® Data Center GPU Max Series 系统上不支持共享系统分配。 然而，通过分配 API 进行共享分配是被支持的。
下表总结了各种类型的内存分配的特性
+——————–+——————+————+————-+ |
Type of allocation | Initial location | Accessible | Accessible | | | | on host? | on device? | +====================+==================+============+=============+ | Host | Host | Yes | Yes | +——————–+——————+————+————-+ | Device | Device | No | Yes | +——————–+——————+————+————-+ | Shared | Host, Device, | Yes | Yes | | | or Unspecified | | | +——————–+——————+————+————-+ | Shared-System | Host | Yes | Yes | | | | | | +——————–+——————+————+————-+
主机分配提供了广泛的可访问性（可以直接从主机和所有支持的设备访问）， 但每次访问的成本可能较高，因为数据通常通过如 PCI Express* 之类的总线发送。
共享分配也提供了广泛的可访问性，但每次访问的成本可能低于主机分配，因为数据会迁移到访问设备。
设备分配有访问限制（不能直接从主机或其他支持的设备访问），但提供了更高的性能，因为访问是对设备本地内存进行的。
OpenMP Runtime 内存分配例程
Intel 编译器支持一些OpenMP运行时例程来执行内存分配。这些例程在下表中显示。
| OpenMP memory allocation routine | Intel extension? | Type of allocation |
|---|---|---|
|
| No | Device |
|
| Yes | Device |
|
| Yes | Host |
|
| Yes | Shared |
请注意，三个例程 omp_target_alloc_device, omp_target_alloc_host 和 omp_target_alloc_shared 是对 OpenMP 规范的 Intel 扩展。
以下示例使用了上述 OpenMP 内存分配例程。将这些与使用 map 子句的示例进行比较。
有关内存分配的更多信息，请参阅：
-
本指南的 SYCL 部分
使用 map 子句
第一个示例使用 map 子句在设备上分配内存并在主机和设备之间拷贝数据。
在以下示例中，通过调用 C/C++ 标准库例程， malloc 在系统内存中分配数组 A, B 和 C。
在第 58 行的 target 构造是在设备上计算数组 C 值的主 kernel。由于需要在计算之前将 C 的值从主机传输到设备， 并在计算结束时从设备传输到主机，因此在此 target 构造上指定了 map(tofrom:C[0:length) 子句。 对于数组 A 和 B ，指定了 map(to: A[0:length], B[0:length])，因为需要将这些数组的值从主机传输到设备， 而设备只读取这些值。在底层， map 子句导致在设备上为数组分配存储空间，并将数据从主机拷贝到设备，反之亦然。
1#include <stdio.h>
2#include <stdlib.h>
3#include <stdint.h>
4#include <math.h>
5#include <omp.h>
6
7#define iterations 100
8#define length 64*1024*1024
9
10int main(void)
11{
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 nstream_time = 0.0;
18
19 // Allocate arrays on the host using plain malloc()
20
21 A = (double *) malloc(bytes);
22 if (A == NULL){
23 printf(" ERROR: Cannot allocate space for A using plain malloc().\n");
24 exit(1);
25 }
26
27 B = (double *) malloc(bytes);
28 if (B == NULL){
29 printf(" ERROR: Cannot allocate space for B using plain malloc().\n");
30 exit(1);
31 }
32
33 C = (double *) malloc(bytes);
34 if (C == NULL){
35 printf(" ERROR: Cannot allocate space for C using plain malloc().\n");
36 exit(1);
37 }
38
39 // Initialize the arrays
40
41 #pragma omp parallel for
42 for (size_t i=0; i<length; i++) {
43 A[i] = 2.0;
44 B[i] = 2.0;
45 C[i] = 0.0;
46 }
47
48 // Perform the computation
49
50 nstream_time = omp_get_wtime();
51 for (int iter = 0; iter<iterations; iter++) {
52 #pragma omp target teams distribute parallel for \
53 map(to: A[0:length], B[0:length]) \
54 map(tofrom: C[0:length])
55 for (size_t i=0; i<length; i++) {
56 C[i] += A[i] + scalar * B[i];
57 }
58 }
59 nstream_time = omp_get_wtime() - nstream_time;
60
61 // Validate and output results
62
63 double ar = 2.0;
64 double br = 2.0;
65 double cr = 0.0;
66 for (int iter = 0; iter<iterations; iter++) {
67 for (int i=0; i<length; i++) {
68 cr += ar + scalar * br;
69 }
70 }
71
72 double asum = 0.0;
73 #pragma omp parallel for reduction(+:asum)
74 for (size_t i=0; i<length; i++) {
75 asum += fabs(C[i]);
76 }
77
78 free(A);
79 free(B);
80 free(C);
81
82 double epsilon=1.e-8;
83 if (fabs(cr - asum)/asum > epsilon) {
84 printf("Failed Validation on output array\n"
85 " Expected checksum: %lf\n"
86 " Observed checksum: %lf\n"
87 "ERROR: solution did not validate\n", cr, asum);
88 return 1;
89 } else {
90 printf("Solution validates\n");
91 double avgtime = nstream_time/iterations;
92 printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
93 }
94
95 return 0;
96}
编译命令:
icpx -fiopenmp -fopenmp-targets=spir64 test_target_map.cpp
运行命令:
OMP_TARGET_OFFLOAD=MANDATORY ZE_AFFINITY_MASK=0.0 LIBOMPTARGET_DEBUG=1 ./a.out
在 iterations 循环内的 target 构造上的 map 子句导致数据（ A, B, C 的值） 在每个 target 区域开始时从主机传输到设备，并导致数据（ C 的值）在每个 target 区域结束时从设备传输到主机。 这些数据传输产生了显著的性能开销。使用 map 子句的更好方法是将整个 iterations 循环放入带有 map 子句的 target data 构造中。 这导致在 iterations 循环开始时发生一次传输，在 iterations 循环结束时再发生一次传输。 下面显示了使用 target data 和 map 子句的修改后的示例。
1#include <stdio.h>
2#include <stdlib.h>
3#include <stdint.h>
4#include <math.h>
5#include <omp.h>
6
7#define iterations 100
8#define length 64*1024*1024
9
10int main(void)
11{
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 nstream_time = 0.0;
18
19 // Allocate arrays on the host using plain malloc()
20
21 A = (double *) malloc(bytes);
22 if (A == NULL){
23 printf(" ERROR: Cannot allocate space for A using plain malloc().\n");
24 exit(1);
25 }
26
27 B = (double *) malloc(bytes);
28 if (B == NULL){
29 printf(" ERROR: Cannot allocate space for B using plain malloc().\n");
30 exit(1);
31 }
32
33 C = (double *) malloc(bytes);
34 if (C == NULL){
35 printf(" ERROR: Cannot allocate space for C using plain malloc().\n");
36 exit(1);
37 }
38
39 // Initialize the arrays
40
41 #pragma omp parallel for
42 for (size_t i=0; i<length; i++) {
43 A[i] = 2.0;
44 B[i] = 2.0;
45 C[i] = 0.0;
46 }
47
48 // Perform the computation
49
50 nstream_time = omp_get_wtime();
51 #pragma omp target data map(to: A[0:length], B[0:length]) \
52 map(tofrom: C[0:length])
53 {
54 for (int iter = 0; iter<iterations; iter++) {
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 }
61 nstream_time = omp_get_wtime() - nstream_time;
62
63 // Validate and output results
64
65 double ar = 2.0;
66 double br = 2.0;
67 double cr = 0.0;
68 for (int iter = 0; iter<iterations; iter++) {
69 for (int i=0; i<length; i++) {
70 cr += ar + scalar * br;
71 }
72 }
73
74 double asum = 0.0;
75 #pragma omp parallel for reduction(+:asum)
76 for (size_t i=0; i<length; i++) {
77 asum += fabs(C[i]);
78 }
79
80 free(A);
81 free(B);
82 free(C);
83
84 double epsilon=1.e-8;
85 if (fabs(cr - asum)/asum > epsilon) {
86 printf("Failed Validation on output array\n"
87 " Expected checksum: %lf\n"
88 " Observed checksum: %lf\n"
89 "ERROR: solution did not validate\n", cr, asum);
90 return 1;
91 } else {
92 printf("Solution validates\n");
93 double avgtime = nstream_time/iterations;
94 printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
95 }
96
97 return 0;
98}
omp_target_alloc
接下来，将上面的示例修改为使用设备分配而不是 map 子句。通过调用 OpenMP runtime 例程 omp_target_alloc 直接 在设备上为数组 A, B 和 C 分配存储空间。该例程接受两个参数：要在设备上分配的字节数，以及要分配存储空间的设备编号。 该例程返回一个设备指针，该指针引用了在设备上分配的存储空间的设备地址。 如果调用 omp_target_alloc 返回 NULL ，则表示分配未成功。
T要在一个 target 构造中访问已分配的内存，可以在该构造上列出一个调用 omp_target_alloc 返回的设备指针，并使用一个 is_device_ptr 子句。 这确保了 kernel 执行前后没有数据传输，因为 kernel 操作的是已经在设备上的数据。
在程序结束时，运行时例程 omp_target_free 用于在设备上释放数组 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 iterations 100
8#define length 64*1024*1024
9
10int main(void)
11{
12 int device_id = omp_get_default_device();
13 size_t bytes = length*sizeof(double);
14 double * __restrict A;
15 double * __restrict B;
16 double * __restrict C;
17 double scalar = 3.0;
18 double nstream_time = 0.0;
19
20 // Allocate arrays in device memory
21
22 A = (double *) omp_target_alloc(bytes, device_id);
23 if (A == NULL){
24 printf(" ERROR: Cannot allocate space for A using omp_target_alloc().\n");
25 exit(1);
26 }
27
28 B = (double *) omp_target_alloc(bytes, device_id);
29 if (B == NULL){
30 printf(" ERROR: Cannot allocate space for B using omp_target_alloc().\n");
31 exit(1);
32 }
33
34 C = (double *) omp_target_alloc(bytes, device_id);
35 if (C == NULL){
36 printf(" ERROR: Cannot allocate space for C using omp_target_alloc().\n");
37 exit(1);
38 }
39
40 // Initialize the arrays
41
42 #pragma omp target teams distribute parallel for \
43 is_device_ptr(A,B,C)
44 for (size_t i=0; i<length; i++) {
45 A[i] = 2.0;
46 B[i] = 2.0;
47 C[i] = 0.0;
48 }
49
50 // Perform the computation 'iterations' number of times
51
52 nstream_time = omp_get_wtime();
53 for (int iter = 0; iter<iterations; iter++) {
54 #pragma omp target teams distribute parallel for \
55 is_device_ptr(A,B,C)
56 for (size_t i=0; i<length; i++) {
57 C[i] += A[i] + scalar * B[i];
58 }
59 }
60 nstream_time = omp_get_wtime() - nstream_time;
61
62 // Validate and output results
63
64 double ar = 2.0;
65 double br = 2.0;
66 double cr = 0.0;
67 for (int iter = 0; iter<iterations; iter++) {
68 for (int i=0; i<length; i++) {
69 cr += ar + scalar * br;
70 }
71 }
72
73 double asum = 0.0;
74 #pragma omp target teams distribute parallel for reduction(+:asum) \
75 map(tofrom: asum) is_device_ptr(C)
76 for (size_t i=0; i<length; i++) {
77 asum += fabs(C[i]);
78 }
79
80 omp_target_free(A, device_id);
81 omp_target_free(B, device_id);
82 omp_target_free(C, device_id);
83
84 double epsilon=1.e-8;
85 if (fabs(cr - asum)/asum > epsilon) {
86 printf("Failed Validation on output array\n"
87 " Expected checksum: %lf\n"
88 " Observed checksum: %lf\n"
89 "ERROR: solution did not validate\n", cr, asum);
90 return 1;
91 } else {
92 printf("Solution validates\n");
93 double avgtime = nstream_time/iterations;
94 printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
95 }
96
97 return 0;
98}
注意事项:
在调用
omp_target_alloc时，指定的设备编号必须是支持的设备之一，而不能是主机设备。这将是分配存储空间的设备。由于数组
A,B和C不可从主机访问，因此需要在 OpenMPtarget区域内完成数组的初始化、kernel 执行和C元素的求和。设备分配只能由
omp_target_alloc调用中指定的设备访问，但可以通过调用omp_target_memcpy将其拷贝到主机或其他设备上分配的内存中。
omp_target_alloc_device
Intel 扩展 omp_target_alloc_device 与 omp_target_alloc 相似。它也是用两个参数调用的： 在设备上分配的字节数，以及要在其上分配存储的设备的编号。该例程返回一个设备指针， 该指针引用设备上分配的存储的设备地址。如果对 omp_target_alloc_device 的调用返回 NULL ，则表示分配不成功。
上述的 omp_target_alloc 示例可以通过简单地将对 omp_target_alloc 的调用替换为对 omp_targer_alloc_device 的调用来重写，如下所示。
在程序结束时， runtime 例程 omp_target_free 用于在设备上释放数组 A, B 和 C 的存储空间。
26 // Allocate arrays in device memory
27
28 A = (double *) omp_target_alloc_device(bytes, device_id);
29 if (A == NULL){
30 printf(" ERROR: Cannot allocate space for A using omp_target_alloc_device().\n");
31 exit(1);
32 }
33
34 B = (double *) omp_target_alloc_device(bytes, device_id);
35 if (B == NULL){
36 printf(" ERROR: Cannot allocate space for B using omp_target_alloc_device().\n");
37 exit(1);
38 }
39
40 C = (double *) omp_target_alloc_device(bytes, device_id);
41 if (C == NULL){
42 printf(" ERROR: Cannot allocate space for C using omp_target_alloc_device().\n");
43 exit(1);
44 }
注意事项:
-
所有适用于omp_target_alloc的上述注意事项也适用于
omp_target_alloc_device。
omp_target_alloc_host
上述示例也可以通过对 A, B 和 C 进行主机分配来重写。这允许内存可以被主机和所有支持的设备访问。
在以下修改后的示例中，调用了 runtime 例程 omp_target_alloc_host （一个 Intel 扩展）来为数组 A, B 和 C 分配存储空间。 该例程接受两个参数：要分配的字节数，以及设备编号。设备编号必须是支持的设备之一，而不能是主机设备。 该例程返回一个指向主机内存中存储位置的指针。如果调用 omp_target_alloc_host 返回 NULL ，这表示分配未成功。
注意在程序顶部指定了指令 requires unified_address。这要求实现保证所有通过 OpenMP API 例程 和指令可访问的设备使用统一地址空间。在此地址空间中，从所有设备看到的指针将始终引用内存中的同一位置， 而且不需要使用 is_device_ptr 子句从设备指针获取设备地址以便在目标区域内使用。当使用 Intel 编译器时， 实际上并不需要 requires unified_address 指令，因为默认情况下就保证了统一地址空间。然而，为了可移植性， 代码包含了该指令。
通过调用 omp_target_alloc_host 返回的指针可以从主机和所有支持的设备访问已分配的内存。由于使用了统一地址空间， 因此在设备上访问内存时不需要在 target 构造上使用任何 map 子句和任何 is_device_ptr 子句。
在程序结束时， runtime 例程 omp_target_free 用于释放数组 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#pragma omp requires unified_address
8
9#define iterations 100
10#define length 64*1024*1024
11
12int main(void)
13{
14 int device_id = omp_get_default_device();
15 size_t bytes = length*sizeof(double);
16 double * __restrict A;
17 double * __restrict B;
18 double * __restrict C;
19 double scalar = 3.0;
20 double nstream_time = 0.0;
21
22 // Allocate arrays in host memory
23
24 A = (double *) omp_target_alloc_host(bytes, device_id);
25 if (A == NULL){
26 printf(" ERROR: Cannot allocate space for A using omp_target_alloc_host().\n");
27 exit(1);
28 }
29
30 B = (double *) omp_target_alloc_host(bytes, device_id);
31 if (B == NULL){
32 printf(" ERROR: Cannot allocate space for B using omp_target_alloc_host().\n");
33 exit(1);
34 }
35
36 C = (double *) omp_target_alloc_host(bytes, device_id);
37 if (C == NULL){
38 printf(" ERROR: Cannot allocate space for C using omp_target_alloc_host().\n");
39 exit(1);
40 }
41
42 // Initialize the arrays
43
44 #pragma omp parallel for
45 for (size_t i=0; i<length; i++) {
46 A[i] = 2.0;
47 B[i] = 2.0;
48 C[i] = 0.0;
49 }
50
51 // Perform the computation
52
53 nstream_time = omp_get_wtime();
54 for (int iter = 0; iter<iterations; iter++) {
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 nstream_time = omp_get_wtime() - nstream_time;
61
62 // Validate and output results
63
64 double ar = 2.0;
65 double br = 2.0;
66 double cr = 0.0;
67 for (int iter = 0; iter<iterations; iter++) {
68 for (int i=0; i<length; i++) {
69 cr += ar + scalar * br;
70 }
71 }
72
73 double asum = 0.0;
74 #pragma omp parallel for reduction(+:asum)
75 for (size_t i=0; i<length; i++) {
76 asum += fabs(C[i]);
77 }
78
79 omp_target_free(A, device_id);
80 omp_target_free(B, device_id);
81 omp_target_free(C, device_id);
82
83 double epsilon=1.e-8;
84 if (fabs(cr - asum)/asum > epsilon) {
85 printf("Failed Validation on output array\n"
86 " Expected checksum: %lf\n"
87 " Observed checksum: %lf\n"
88 "ERROR: solution did not validate\n", cr, asum);
89 return 1;
90 } else {
91 printf("Solution validates\n");
92 double avgtime = nstream_time/iterations;
93 printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
94 }
95
96 return 0;
97}
注意事项:
当调用
omp_target_alloc_host时，指定的设备号必须是支持的设备之一，不能是主机设备。由于数组
A、B``和 ``C可以从主机和设备访问，因此可以在主机（在target结构体外部）或 设备（在target结构体内部）上初始化数组并对C的元素进行求和。Intel® Data Center GPU Max 系列不支持对主机分配（即通过
omp_target_alloc_host分配内存） 进行原子操作（或使用原子操作的算法，如某些 reduction ）。请在通过omp_target_alloc_device分配的内存上使用原子操作。
omp_target_alloc_shared
上述示例已修改，以使用共享分配代替主机分配。runtime 例程 omp_target_alloc_shared 被调用来为数组 A、 B``和 ``C 分配存储空间。 该例程接受两个参数：在设备上分配的字节数和设备号。设备号必须是支持的设备之一，不能是主机设备。 该例程返回一个指向共享内存中存储位置的指针。如果调用 omp_target_alloc_shared 返回 NULL ，则表示分配未成功。
注意，在程序顶部指定了 requires unified_addres 指令，以确保可移植性。
通过调用 omp_target_alloc_share 返回的指针可以从主机和所有支持的设备访问存储。由于使用了统一地址空间， 因此在 targe 结构体上访问设备的内存时，不需要 map 子句和 is_device_pt 子句。
在程序结束时， runtime 例 omp_target_fre 用于为 A、 B 和 C 释放存储空间。
28 // Allocate arrays in shared memory
29
30 A = (double *) omp_target_alloc_shared(bytes, device_id);
31 if (A == NULL){
32 printf(" ERROR: Cannot allocate space for A using omp_target_alloc_shared().\n");
33 exit(1);
34 }
35
36 B = (double *) omp_target_alloc_shared(bytes, device_id);
37 if (B == NULL){
38 printf(" ERROR: Cannot allocate space for B using omp_target_alloc_shared().\n");
39 exit(1);
40 }
41
42 C = (double *) omp_target_alloc_shared(bytes, device_id);
43 if (C == NULL){
44 printf(" ERROR: Cannot allocate space for C using omp_target_alloc_shared().\n");
45 exit(1);
46 }
注意事项:
当调用
omp_target_alloc_shared时，指定的设备号必须是支持的设备之一，不能是主机设备。由于数组可以从主机和设备访问，因此初始化和验证可以在主机或设备（在
target结构体内部）上完成。不支持从主机和设备同时访问通过
omp_target_alloc_shared分配的内存。
omp_target_memcpy
以下示例展示了如何使用 runtime 例程 omp_target_memcpy 将内存从主机拷贝到设备，以及从设备拷贝到主机。 首先使用普通的 malloc 在系统内存中为数组 h_A、 h_B 和 h_C 分配空间，并进行初始化。 然后使用 omp_target_alloc 在设备上为对应的数组 d_A 、 d_B 和 d_C 分配空间。
在第 104 行的 target 结构体开始之前，通过调用 omp_target_memcpy，将 h_A、 h_B 和 h_C 中的值拷贝到 d_A 、 d_B 和 d_C 。在 target 区域之后， 在设备上计算出的新的 d_C 值通过调用 omp_target_memcpy 拷贝到 h_C。
1#include <stdio.h>
2#include <stdlib.h>
3#include <stdint.h>
4#include <math.h>
5#include <omp.h>
6
7#define iterations 100
8#define length 64*1024*1024
9
10int main(void)
11{
12 int device_id = omp_get_default_device();
13 int host_id = omp_get_initial_device();
14 size_t bytes = length*sizeof(double);
15 double * __restrict h_A;
16 double * __restrict h_B;
17 double * __restrict h_C;
18 double * __restrict d_A;
19 double * __restrict d_B;
20 double * __restrict d_C;
21 double scalar = 3.0;
22 double nstream_time = 0.0;
23
24 // Allocate arrays h_A, h_B, and h_C on the host using plain malloc()
25
26 h_A = (double *) malloc(bytes);
27 if (h_A == NULL){
28 printf(" ERROR: Cannot allocate space for h_A using plain malloc().\n");
29 exit(1);
30 }
31
32 h_B = (double *) malloc(bytes);
33 if (h_B == NULL){
34 printf(" ERROR: Cannot allocate space for h_B using plain malloc().\n");
35 exit(1);
36 }
37
38 h_C = (double *) malloc(bytes);
39 if (h_C == NULL){
40 printf(" ERROR: Cannot allocate space for h_C using plain malloc().\n");
41 exit(1);
42 }
43
44 // Allocate arrays d_A, d_B, and d_C on the device using omp_target_alloc()
45
46 d_A = (double *) omp_target_alloc(bytes, device_id);
47 if (d_A == NULL){
48 printf(" ERROR: Cannot allocate space for d_A using omp_target_alloc().\n");
49 exit(1);
50 }
51
52 d_B = (double *) omp_target_alloc(bytes, device_id);
53 if (d_B == NULL){
54 printf(" ERROR: Cannot allocate space for d_B using omp_target_alloc().\n");
55 exit(1);
56 }
57
58 d_C = (double *) omp_target_alloc(bytes, device_id);
59 if (d_C == NULL){
60 printf(" ERROR: Cannot allocate space for d_C using omp_target_alloc().\n");
61 exit(1);
62 }
63
64 // Initialize the arrays on the host
65
66 #pragma omp parallel for
67 for (size_t i=0; i<length; i++) {
68 h_A[i] = 2.0;
69 h_B[i] = 2.0;
70 h_C[i] = 0.0;
71 }
72
73 // Call omp_target_memcpy() to copy values from host to device
74
75 int rc = 0;
76 rc = omp_target_memcpy(d_A, h_A, bytes, 0, 0, device_id, host_id);
77 if (rc) {
78 printf("ERROR: omp_target_memcpy(A) returned %d\n", rc);
79 exit(1);
80 }
81
82 rc = omp_target_memcpy(d_B, h_B, bytes, 0, 0, device_id, host_id);
83 if (rc) {
84 printf("ERROR: omp_target_memcpy(B) returned %d\n", rc);
85 exit(1);
86 }
87
88 rc = omp_target_memcpy(d_C, h_C, bytes, 0, 0, device_id, host_id);
89 if (rc) {
90 printf("ERROR: omp_target_memcpy(C) returned %d\n", rc);
91 exit(1);
92 }
93
94 // Perform the computation
95
96 nstream_time = omp_get_wtime();
97 for (int iter = 0; iter<iterations; iter++) {
98 #pragma omp target teams distribute parallel for \
99 is_device_ptr(d_A,d_B,d_C)
100 for (size_t i=0; i<length; i++) {
101 d_C[i] += d_A[i] + scalar * d_B[i];
102 }
103 }
104 nstream_time = omp_get_wtime() - nstream_time;
105
106 // Call omp_target_memcpy() to copy values from device to host
107
108 rc = omp_target_memcpy(h_C, d_C, bytes, 0, 0, host_id, device_id);
109 if (rc) {
110 printf("ERROR: omp_target_memcpy(A) returned %d\n", rc);
111 exit(1);
112 }
113
114 // Validate and output results
115
116 double ar = 2.0;
117 double br = 2.0;
118 double cr = 0.0;
119 for (int iter = 0; iter<iterations; iter++) {
120 for (int i=0; i<length; i++) {
121 cr += ar + scalar * br;
122 }
123 }
124
125 double asum = 0.0;
126 #pragma omp parallel for reduction(+:asum)
127 for (size_t i=0; i<length; i++) {
128 asum += fabs(h_C[i]);
129 }
130
131 free(h_A);
132 free(h_B);
133 free(h_C);
134 omp_target_free(d_A, device_id);
135 omp_target_free(d_B, device_id);
136 omp_target_free(d_C, device_id);
137
138 double epsilon=1.e-8;
139 if (fabs(cr - asum)/asum > epsilon) {
140 printf("Failed Validation on output array\n"
141 " Expected checksum: %lf\n"
142 " Observed checksum: %lf\n"
143 "ERROR: solution did not validate\n", cr, asum);
144 return 1;
145 } else {
146 printf("Solution validates\n");
147 double avgtime = nstream_time/iterations;
148 printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
149 }
150
151 return 0;
152}
性能考虑因素
在上述示例中（使用 map 子句， omp_target_alloc ， omp_target_alloc_device ， omp_target_alloc_host ， omp_target_alloc_shared ， omp_target_memcpy ）， 主 kernel 是计算数组 C 值的 target 结构。为了获得更准确的时间，这个 target 结构被包含在一个循环中，所以部署发生了 iterations 次（其中 iterations = 100 ）。 通过将循环 iterations 所用的总时间除以 100 ，计算出 kernel 所用的平均时间。
56 // Perform the computation 'iterations' number of times
57
58 nstream_time = omp_get_wtime();
59 for (int iter = 0; iter<iterations; iter++) {
60 #pragma omp target teams distribute parallel for \
61 is_device_ptr(A,B,C)
62 for (size_t i=0; i<length; i++) {
63 C[i] += A[i] + scalar * B[i];
64 }
65 }
66 nstream_time = omp_get_wtime() - nstream_time;
LIBOMPTARGET_DEBUG=1 输出显示，所有上述示例都具有相同的 ND_range 分区。
Target LEVEL0 RTL --> Allocated a device memory 0xff00000020200000 Libomptarget --> omp_target_alloc returns device ptr 0xff00000020200000 Libomptarget --> Call to omp_target_alloc for device 0 requesting 536870912 bytes Libomptarget --> Call to omp_get_num_devices returning 1 Libomptarget --> Call to omp_get_initial_device returning 1 Libomptarget --> Checking whether device 0 is ready. Libomptarget --> Is the device 0 (local ID 0) initialized? 1 Libomptarget --> Device 0 is ready to use.
下表显示了在使用的特定 GPU （仅 1 堆栈）上运行各个版本时， kernel 所用的平均时间。
| Version | Time (seconds) |
|---|---|
|
| 0.183604 |
|
| 0.012757 |
|
| 0.002501 |
|
| 0.002499 |
|
| 0.074412 |
|
| 0.012491 |
|
| 0.011072 |
上述性能数据显示，使用 map 子句的版本是最慢的版本（ 0.183604秒 ）。这是因为在每次 kernel 启动的开始和结束时都会发生数据传输。 主 kernel 启动了 100 次。在每次 kernel 启动的开始，设备上为数组 A、 B``和 ``C 分配了存储空间，并将这些数组的值从主机拷贝到设备。 在 kernel 结束时，将数组 C 的值从设备拷贝到主机。将整个 iterations 循环放在带有 map 子句的 target data 结构中，可以将运行时间减少到 0.012757 秒， 因为传输只在 iterations 循环中的第一个 kernel 启动时发生一次，然后在该循环中的最后一个 kernel 之后再发生一次。
omp_target_alloc 和 omp_target_alloc_device 版本具有最佳性能（分别为 0.002501 秒和 0.002499 秒）。在这些版本中， 直接在设备内存中为 A、 B 和 C 分配存储空间，因此设备上的访问是从设备本地内存进行的。 这对于在设备端使用临时数组的应用程序来说是一个有用的模型。这些数组永远不需要在主机上访问。 在这种情况下，建议将临时数组分配到设备上，并不用担心数据传输，如本示例所示。
使用 omp_target_alloc_shared 版本也表现良好，但稍慢（ 0.012491 秒）。在此版本中，为 A、 B 和 C 分配了共享内存中的存储空间。 因此，数据可以在主机和设备之间迁移。迁移有一定的开销，但迁移后，在设备上的访问会从更快的设备本地内存进行。 在此版本中，数组的初始化发生在主机上。在第一次 kernel 启动时，数组被迁移到设备上，并且 kernel 在设备上本地访问数组。 最后，在主机执行 reduction 计算之前，整个 C 数组被迁回到主机。
使用 omp_target_alloc_host 版本（ 0.074412 秒）所需时间几乎是 omp_target_alloc_shared 版本所需时间的 6 倍。 这是因为在主机内存中分配的数据不会从主机迁移到设备。当 kernel 试图访问数据时，数据通常通过连接设备和主机的总线（如 PCI Express ）发送。 这比访问本地设备内存要慢。如果设备只偶尔访问数组的一小部分，则可以使用 omp_target_alloc_host 将该数组分配到主机内存中。 然而，如果设备端频繁访问数组，则应将其保留在设备内存中。将数据保留在主机内存中并通过 PCI 访问它会降低性能。
最后，关于数据传输的注意事项：可以通过搜索 "Libomptarget --> Moving" 来查看 map 版本中传输的数据量。注意每次启动主 kernel 都会产生以下数据传输:
$ grep "Libomptarget --> Moving" test_target_map.debug Libomptarget --> Moving 536870912 bytes (hst:0x00007f1a5fc8b010) -> (tgt:0xff00000000200000) Libomptarget --> Moving 536870912 bytes (hst:0x00007f1a9fc8d010) -> (tgt:0xff00000020200000) Libomptarget --> Moving 536870912 bytes (hst:0x00007f1a7fc8c010) -> (tgt:0xff00000040200000) Libomptarget --> Moving 536870912 bytes (tgt:0xff00000000200000) -> (hst:0x00007f1a5fc8b010)
另一方面， omp_target_alloc_... 版本中的数据传输由 runtime 系统的较低层处理。因此，在 LIBOMPTARGET_DEBUG=1 输出中搜索 "Libomptarget --> Moving" 将不会显示发生的数据传输。
Fortran 示例
下面展示了使用 target data 和 map 子句的 Fortran 示例。
1 program main 2 use iso_fortran_env 3 use omp_lib 4 implicit none 5 6 integer, parameter :: iterations=100 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 :: err, i, iter 16 17 ! 18 ! Allocate arrays on the host using plain allocate 19 20 allocate( A(length), stat=err ) 21 if (err .ne. 0) then 22 print *, "Allocation of A returned ", err 23 stop 1 24 endif 25 26 allocate( B(length), stat=err ) 27 if (err .ne. 0) then 28 print *, "Allocation of B returned ", err 29 stop 1 30 endif 31 32 allocate( C(length), stat=err ) 33 if (err .ne. 0) then 34 print *, "Allocation of C returned ", err 35 stop 1 36 endif 37 38 ! 39 ! Initialize the arrays 40 41 !$omp parallel do 42 do i = 1, length 43 A(i) = 2.0 44 B(i) = 2.0 45 C(i) = 0.0 46 end do 47 48 ! 49 ! Perform the computation 50 51 nstream_time = omp_get_wtime() 52 !$omp target data map(to: A, B) map(tofrom: C) 53 54 do iter = 1, iterations 55 !$omp target teams distribute parallel do 56 do i = 1, length 57 C(i) = C(i) + A(i) + scalar * B(i) 58 end do 59 end do 60 61 !$omp end target data 62 nstream_time = omp_get_wtime() - nstream_time 63 64 ! 65 ! Validate and output results 66 67 ar = 2.0 68 br = 2.0 69 cr = 0.0 70 do iter = 1, iterations 71 do i = 1, length 72 cr = cr + ar + scalar * br 73 end do 74 end do 75 76 asum = 0.0 77 !$omp parallel do reduction(+:asum) 78 do i = 1, length 79 asum = asum + abs(C(i)) 80 end do 81 82 if (abs(cr - asum)/asum > epsilon) then 83 write(*,110) "Failed Validation on output array: Expected =", cr, ", Observed =", asum 84 else 85 avgtime = nstream_time/iterations 86 write(*,120) "Solution validates: Checksum =", asum, ", Avg time (s) =", avgtime 87 endif 88 89110 format (A, F20.6, A, F20.6) 90120 format (A, F20.6, A, F10.6) 91 92 deallocate(A) 93 deallocate(B) 94 deallocate(C) 95 96 end program main
下面展示了使用 omp_target_alloc_device 的 Fortran 示例。在此示例中，使用分配器 omp_target_device_mem_alloc 的 allocate 指令在设备上为数组 A、 B 和 C 分配空间。 在 target data 指令（第 37 行）上使用了 use_device_addr(A, B, C) 子句，以指示数组具有设备地址，并且应在 target 区域中使用这些地址。
1 use iso_fortran_env 2 use omp_lib 3 implicit none 4 5 integer, parameter :: iterations=100 6 integer, parameter :: length=64*1024*1024 7 real(kind=REAL64), parameter :: epsilon=1.D-8 8 real(kind=REAL64), allocatable :: A(:) 9 real(kind=REAL64), allocatable :: B(:) 10 real(kind=REAL64), allocatable :: C(:) 11 real(kind=REAL64) :: scalar=3.0 12 real(kind=REAL64) :: ar, br, cr, asum 13 real(kind=REAL64) :: nstream_time, avgtime 14 integer :: i, iter 15 16 ! 17 ! Allocate arrays in device memory 18 19 !$omp allocate allocator(omp_target_device_mem_alloc) 20 allocate(A(length)) 21 22 !$omp allocate allocator(omp_target_device_mem_alloc) 23 allocate(B(length)) 24 25 !$omp allocate allocator(omp_target_device_mem_alloc) 26 allocate(C(length)) 27 28 ! 29 ! Begin target data 30 31 !$omp target data use_device_addr(A, B, C) 32 33 ! 34 ! Initialize the arrays 35 36 !$omp target teams distribute parallel do 37 do i = 1, length 38 A(i) = 2.0 39 B(i) = 2.0 40 C(i) = 0.0 41 end do 42 43 ! 44 ! Perform the computation 45 46 nstream_time = omp_get_wtime() 47 do iter = 1, iterations 48 !$omp target teams distribute parallel do 49 do i = 1, length 50 C(i) = C(i) + A(i) + scalar * B(i) 51 end do 52 end do 53 nstream_time = omp_get_wtime() - nstream_time 54 55 ! 56 ! Validate and output results 57 58 ar = 2.0 59 br = 2.0 60 cr = 0.0 61 do iter = 1, iterations 62 do i = 1, length 63 cr = cr + ar + scalar * br 64 end do 65 end do 66 67 asum = 0.0 68 !$omp target teams distribute parallel do reduction(+:asum) & 69 !$omp map(tofrom: asum) 70 do i = 1, length 71 asum = asum + abs(C(i)) 72 end do 73 74 ! 75 ! End target data 76 77 !$omp end target data 78 79 if (abs(cr - asum)/asum > epsilon) then 80 write(*,110) "Failed Validation on output array: Expected =", cr, ", Observed =", asum 81 else 82 avgtime = nstream_time/iterations 83 write(*,120) "Solution validates: Checksum =", asum, ", Avg time (s) =", avgtime 84 endif 85 86110 format (A, F20.6, A, F20.6) 87120 format (A, F20.6, A, F10.6) 88 89 deallocate(A) 90 deallocate(B) 91 deallocate(C) 92 93 end program main
扫一扫
788
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
