使用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指令在一條指令中計算（AB）+C。FMA結果與乘法后加指令的結果之間的區別在于FMA中間產物aB在加到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中。

總結

以上是生活随笔為你收集整理的使用PCAST检测散度以比较GPU和CPU结果的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。