使用PCAST检测散度以比较GPU和CPU结果

并行编译器辅助软件测试（PCAST）是英伟达HPC FORTRAN、C++和C编译器中的一个特性。PCAST有两个用例。一个新的处理器或新的编译程序的部分或新的时间标志首先被编译。您可能需要测试新库是否会产生相同的结果，或者测试添加OpenMP并行、启用自动矢量化（-Mvect=simd）或从X86系统移植到OpenPOWER或Arm的安全性。这个用例的工作原理是在需要比较中间结果的地方向应用程序添加pcast_compare调用或compare指令。在初始运行时，这些结果保存在一个golden文件中，知道结果是正确的。在测试运行期间，相同的调用或指令将计算的中间结果与保存的结果进行比较，并报告差异。

第二个用例特定于NVIDIA OpenACC实现。这将GPU计算与CPU上运行的相同程序进行比较。在这种情况下，所有的计算构造都是在CPU和GPU上冗余完成的。然后可以将GPU结果与CPU结果和报告的差异进行比较。

PCAST with a golden file

在这个用例中，好的结果被保存到一个golden文件中，并将测试结果与这些结果进行比较。这是通过向程序中添加pcast_compare调用或compare指令来完成的。它由PCAST_COMPARE环境变量控制。下面的代码示例显示了调用slover的过程：

void solve(double* a, double* b, double* r, int* pivot, int n){

int fail;

dgesv(n, 1, a, n, pivot, b, n, &fail);

}

….

solve(a, b, r, pivot, n);

要将结果与NAG版本进行比较：

#include

void solve(double* a, double* b, double* r, int* pivot, int n){

NagError fail;

nag_dgesv(NAG_RowMajor, n, 1, a, n, pivot, b, n, &fail);

}

….

solve(a, b, r, pivot, n);

通过在调用solve之后添加一个、两个或三个PCAST_compare调用或compare指令，可以使用PCAST保存和比较这些结果。在本例中，您可能需要比较b向量中的结果，以及a矩阵和pivot索引向量的LU分解。这很容易通过指令完成：

solve(a, b, r, pivot, n);

#pragma pcast compare(b[0:n])

#pragma pcast compare(a[0:n*n])

#pragma pcast compare(pivot[0:n])

或者，可以将直接调用插入到pcast_compare：

solve(a, b, r, pivot, n);

pcast_compare(b, "double", n, "b", "solve.c", "myfunc", 1);

pcast_compare(a, "double", n*n, "a", solve.c", "myfunc", 2);

pcast_compare(pivot, "int", n, "pivot", solve.c", "myfunc", 3);

• The first argument to pcast_compare is the address of the data to be saved or compared.
• The second argument is a string containing the data type, here double precision.
• The third argument is the number of elements to compare.
• The next three arguments are strings which pcast_compare treats as the variable name, source file name, and function name where the call appears.
• The last argument is an integer, which pcast_compare treats as a line number.

pcast_compare的第一个参数是要保存或比较的数据的地址。

第二个参数是一个包含数据类型的字符串，这里是double precision。

第三个参数是要比较的元素数。

接下来的三个参数是pcast_compare将其视为变量名、源文件名和调用出现的函数名。

最后一个参数是一个整数，pcast_compare将其视为行号。

当然，可以对最后四个参数使用任何字符串或值；

仅用于输出，并确保在测试运行中进行的调用序列与第一次golden运行时相同。

PCAST支持的数据类型包括C、C++和FORTRAN的基本数字类型：

• float, real, real(4)
• double, double precision, real(8)
• complex, float _Complex
• double _Complex, complex(8)
• short
• int, integer, integer(4)
• long, integer(8)
• unsigned short
• unsigned int
• unsigned long

compare指令或pcast_compare调用写入或读取名为pcast的golden数据文件pcast_compare.dat比较数据默认情况下。如果文件不存在，则运行时假定这是第一次golden运行，创建该文件，并用计算的数据填充它。如果该文件存在，则运行时假定这是一个测试运行，读取该文件，并将计算的数据与该文件中保存的数据进行比较。

可以使用PCAST_COMPARE环境变量更改文件名。默认行为是将任何偏差（无论多小）视为错误并报告前50个偏差。PCAST_COMPARE环境变量可用于容忍小偏差，或更改生成的输出。

PCAST with OpenACC and Autocompare

对于OpenACC程序，PCAST包含一个选项，可以简化GPU内核对相应CPU代码的测试。启用后，编译器会为每个计算构造生成CPU和GPU代码。在运行时，CPU和GPU版本都是冗余运行的。CPU代码读取和修改系统内存中的值，GPU读取并修改设备内存中的值。然后，可以在要将GPU计算的值与CPU计算的值进行比较的点插入acc_compare的调用。本质上，这种方法将CPU代码视为计算golden值。它不读写文件，而是计算和比较内存中的值。使用-GPU=redundant启用冗余CPU+GPU代码生成模式。

一个更有趣和半自动的方法是，允许运行时在从设备内存下载值时自动比较它们。这是通过-gpu=autocompare编译器标志启用的，该标志还启用了冗余选项。这将在CPU和GPU上冗余地运行每个计算结构，并比较结果，而不更改程序本身。

下面的代码示例测试矩阵或向量积的结果：

void matvec(double* a, double* x, double* v, int n){

#pragma acc parallel loop copyin(a[0:n*n], x[0:n]) copyout(v[0:n])

for (int i = 0; i < n; ++i) {

double r = 0.0;

#pragma acc loop reduction(+:r)

for (int j = 0; j < n; ++i)

r += a[i*n+j] * x[j];

v[i] = r;

}

}

If you build this program with the -acc=gpu flag, without autocompare, the generated code performs the following sequence of operations:

1. Allocate space for a, x, and v.
2. Copy the input values for a and x from host to device memory.
3. Launch the compute kernel on the device.
4. Copy the output values for v from device memory back to v in host memory.
5. Deallocate space for a, x, and v.

If instead you build with the -acc=gpu:autocompare flag, the sequence of operations is as follows:

1. Allocate space for a, x, and v.
2. Copy the input values for a and x from host to device memory.
3. Launch the compute kernel on the device.
4. Redundantly execute the loop on the CPU.
5. Download the GPU-computed values for v from device memory back to temporary memory.
6. Compare the downloaded values with those computed by the CPU.
7. Deallocate space for a, x, and v.

编译器和OpenACC运行时，说明数据类型，元素的数量来自data子句。结果有点像对下载的数据执行pcast_compare调用。事实上，这些比较可以由同一个PCAST_COMPARE环境变量控制，就像使用golden文件的PCAST一样。有关更多信息，请参阅本文后面的PCAST_COMPARE环境变量部分。

自动比较功能只在下载到系统内存时比较数据。在数据区域（数据已经存在于设备上）中的某些计算构造之后，要比较数据，有三种方法可以在程序中的任何点进行比较。

首先，可以插入一个update self指令来下载要比较的数据。启用autocompare选项后，任何使用update self指令下载的数据都将从GPU下载，并与主机CPU上计算的值进行比较。在本例中，插入以下指令将执行比较：

…

v[i] = r;

}

#pragma acc update host(v[0:n])

}

或者，可以添加对acc_compare的调用，它将GPU上的值与主机内存中的相应值进行比较。acc_compare例程只有两个参数：要比较的数据的地址和要比较的元素的数量。数据类型在OpenACC运行时中可用，因此不需要指定。在本例中，调用如下：

…

v[i] = r;

}

acc_compare(v, n);

}

可以对设备内存中存在的任何变量或数组调用acc_compare。也可以不带参数调用acc_compare_all，将设备内存中的所有值与主机内存中的相应值进行比较。它们只与-gpu=redundant或autocompare一起使用，其中所有的计算构造都在CPU和gpu上执行，并且值应该相同。

最后，您可以使用新的acc compare指令：

PCAST_COMPARE environment variable

PCAST_COMPARE环境变量有几个有用的设置来控制PCAST。可以指定多个以逗号分隔的设置，例如在以下命令中：

export PCAST_COMPARE=rel=5,summary,file=myfile.dat

The PCAST_COMPARE options are listed below. The first three file options below do not apply to OpenACC autocompare or acc_compare.

• file=name—Use this filename when writing and reading a golden file.
• create—Create the golden file, even if it already exists.
• compare—Compare results against the golden file. If the file does not exist, issue an error message.
• report=n—Report up to n differences; the default is 50.
• stop—Stop after finding one data block with differences.
• summary—Print a summary of comparisons at the end of the run.
• abs=n—Tolerate or ignore differences for floating-point data when the absolute difference abs(a-b) is less than 10^(-n).
• rel=n—Tolerate ignore differences for floating-point data when the relative difference abs((a-b)/a) is less than 10^(-n).
• ieee—Report differences in IEEE NaN values as well.

pcast_compare.dat比较数据文件可能很大。对于包含大数据结构写入操作的循环，不需要多次迭代就可以创建数GB的文件。写入和重读此文件也可能相应地缓慢。对于大型应用程序，我们建议您谨慎地使用pcast_compare调用或pcast compare指令，然后在开始出现差异时进入。

比较不是以线程安全的方式进行的，并且比较不考虑哪个线程正在进行比较。添加一个线程时，你应该选择一个线程与一个线程进行比较。当使用golden文件将PCAST添加到MPI程序时，多个MPI列组将写入或读取同一个文件，除非有一个脚本使用PCAST_COMPARE重命名该文件，并与文件名中编码的MPI列组进行比较。

在更改数据类型后，当前没有比较结果的功能。例如，可能想知道在将精度从双精度降低到单精度后，结果是否有显著差异。PCAST无法将golden文件中的双精度值与测试运行中计算的单个精度值进行比较。目前，PCAST无法比较结构体或派生类型，除非它们可以作为数组进行比较。

当使用OpenACC autocompare或带acc_compare的冗余执行时，不能使用CUDA统一内存或-gpu=managed选项。OpenACC-PCAST比较依赖于让GPU在设备内存中执行其计算，独立于在主机内存中执行冗余计算的CPU。

如果有一些计算或数据移动超出了OpenACC的控制范围，只会影响主机或设备内存，那么在使用PCAST时必须考虑到这一点。例如，必须将MPI数据传输到主机内存，并将其复制到设备上。如果使用支持CUDA的MPI并在两个GPU之间直接执行MPI数据传输，则主机内存值已过时，必须更新。类似地，如果您在设备上有任何cuBLAS或cuSolver调用，那么主机内存值就是过时的。这很可能是OpenACC host_data结构的地方，它表示设备内存中的某些内容正在被处理。

Important considerations

即使修改后的程序是正确的，计算结果也可能存在差异，特别是对于浮点数据。如果修改改变了某些内在函数的实现方式，可能会产生差异。当转移到新处理器或使用不同的编译器时，这是非常可能的。

如果修改后的程序使用的融合乘法加法（FMA）指令与原始程序不同，也可能产生差异。FMA指令在一条指令中计算（A*B）+C。FMA结果与乘法后加指令的结果之间的区别在于FMA中间产物a*B在加到C之前不是四舍五入的。事实上，FMA中间结果在加法中携带了更重要的位，因此您可以说，对于better的某些定义，结果更好，但是关键是它是不同的。

对于并行操作，也可能在几个方面产生差异。每次运行并行程序时，原子操作可能以不同的顺序发生。并行约化，特别是并行求和，以与顺序程序不同的顺序累积结果，从而产生不同的舍入误差。

考虑到存在差异的可能性，区分显著差异和无关紧要的差异是很重要的。在修改程序之后，要求所有浮点计算都是位精确的已经不再合理了。只有能确定什么时候差异是显著的，或者差异有多大才有意义。有许多正在进行的工作，以确定和隔离所有这些差异的原因，甚至有专门讨论这一点的研讨会。希望这最终会使这个过程更加自动化，程序员的工作量减少。

本文，OpenACC冗余执行和自动比较有一个已知的限制。带有if子句的compute构造没有在CPU和GPU上以冗余方式正确执行，即使条件为真。

Future directions

我们正在探索压缩生成的数据文件，这将以牺牲一些额外的计算时间来解决文件大小问题。想探索如何并行地进行比较，无论是在CPU上，还是在使用OpenACC时，在GPU上进行比较。还希望允许struct和派生类型。

设计的程序是为了测试频率和控制频率而设计的。也许可以使用PCAST_COMPARE环境变量来指定文件名或函数名，而实际上只在该文件或函数中进行比较。尤其是使用指令接口时，这可以允许用户将指令留在程序中，并通过粗略、不频繁的比较来启动调试过程。如果在调用某个模块后出现差异，请在该模块中启用更频繁的比较，并重复此过程，直到找到差异的原因。

和用户的经验已经确定，PCAST中发现的大多数错误都是由于缺少数据或更新指令造成的，其中CPU或GPU正在处理另一个更新的过时数据值。

当比较CPU和GPU的运行时，经常会有差异源于求和的减少。GPU上的并行代码以不同于CPU的顺序累积和减少，因此舍入误差的累积也不同。正在寻找减少这些差异的方法，以满足程序员的要求，即差异只是由于求和时的舍入误差，而不是别的错误。

Summary

文章描述了支持软件测试的PCAST特性，特别是将一个已知良好的程序的结果与一个假定要执行相同计算的修改程序的测试运行进行比较。修改可以是源代码更改、链接到不同的库或内部版本差异（编译器标志），也可以是更改为新的处理器类型。

PCAST由C、C++和FORTRAN HPC编译器支持，包括在英伟达HPC SDK中。