CUDA 3：錯誤檢測
編寫CUDA程式難免出現錯誤，編譯錯誤這種能在編譯過程中被編譯器捕抓的還好，如果是編譯期間沒有被發現而在執行時出現，這種錯誤更難排查。本文著重討論如何檢測執行時刻的錯誤。
一個檢測CUDA執行時錯誤的宏函式
之前提到的
記憶體分配函式cudaMalloc()、釋放記憶體函式cudaFree()以及傳輸資料的函式cudaMemcpy()，都有一個cudaError_t型別的返回值，代表一種錯誤資訊，只有返回cudaSuccess時才代表成功呼叫了cuda的API函式
。依這樣的原則可以寫一個頭檔案（error。cuh），其包含一個檢測執行時刻錯誤的宏函式（macro function）：
#pragma once
#include
#define CHECK（call） \
do \
{ \
const cudaError_t error_code = call； \
if （error_code ！= cudaSuccess） \
{ \
printf（“CUDA Error：\n”）； \
printf（“ File： %s\n”， __FILE__）； \
printf（“ Line： %d\n”， __LINE__）； \
printf（“ Error code： %d\n”， error_code）； \
printf（“ Error text： %s\n”， \
cudaGetErrorString（error_code））； \
exit（1）； \
} \
} while （0）
這段程式碼中第一行
#pragma once是一個預處理命令
，用於確保當前檔案在一個編譯單元中不被重複包含。宏函式的名為CHECK，引數
call是一個CUDA執行時API函式
；在定義宏時如果一行寫不下則需要
\在行末以表示續行
。此外
cudaGetErrorString（）是一個將錯誤代號轉換為錯誤文字描述的API函式
。使用該宏函式時只要將一個CUDA執行時API當作引數傳入該宏函式即可。
檢查執行時API函式
將上面的錯誤檢測用於陣列相加很簡單，
只需要在指令碼開頭包含上述的標頭檔案並將CUDA執行時API函式都用宏函式CHECK包裝
即可：
#include
“error。cuh”
#include
#include
const
double
EPSILON
=
1。0e-15
；
const
double
a
=
1。23
；
const
double
b
=
2。34
；
const
double
c
=
3。57
；
void
__global__
add
（
const
double
*
x
，
const
double
*
y
，
double
*
z
，
const
int
N
）；
void
check
（
const
double
*
z
，
const
int
N
）；
int
main
（
void
）
{
const
int
N
=
100000000
；
const
int
M
=
sizeof
（
double
）
*
N
；
double
*
h_x
=
（
double
*
）
malloc
（
M
）；
double
*
h_y
=
（
double
*
）
malloc
（
M
）；
double
*
h_z
=
（
double
*
）
malloc
（
M
）；
for
（
int
n
=
0
；
n
<
N
；
++
n
）
{
h_x
［
n
］
=
a
；
h_y
［
n
］
=
b
；
}
double
*
d_x
，
*
d_y
，
*
d_z
；
CHECK
（
cudaMalloc
（（
void
**
）
&
d_x
，
M
））；
CHECK
（
cudaMalloc
（（
void
**
）
&
d_y
，
M
））；
CHECK
（
cudaMalloc
（（
void
**
）
&
d_z
，
M
））；
CHECK
（
cudaMemcpy
（
d_x
，
h_x
，
M
，
cudaMemcpyDeviceToHost
））；
CHECK
（
cudaMemcpy
（
d_y
，
h_y
，
M
，
cudaMemcpyDeviceToHost
））；
const
int
block_size
=
128
；
const
int
grid_size
=
（
N
+
block_size
-
1
）
/
block_size
；
add
<<<
grid_size
，
block_size
>>>
（
d_x
，
d_y
，
d_z
，
N
）；
CHECK
（
cudaMemcpy
（
h_z
，
d_z
，
M
，
cudaMemcpyDeviceToHost
））；
check
（
h_z
，
N
）；
free
（
h_x
）；
free
（
h_y
）；
free
（
h_z
）；
CHECK
（
cudaFree
（
d_x
））；
CHECK
（
cudaFree
（
d_y
））；
CHECK
（
cudaFree
（
d_z
））；
return
0
；
}
void
__global__
add
（
const
double
*
x
，
const
double
*
y
，
double
*
z
，
const
int
N
）
{
const
int
n
=
blockDim
。
x
*
blockIdx
。
x
+
threadIdx
。
x
；
if
（
n
<
N
）
{
z
［
n
］
=
x
［
n
］
+
y
［
n
］；
}
}
void
check
（
const
double
*
z
，
const
int
N
）
{
bool
has_error
=
false
；
for
（
int
n
=
0
；
n
<
N
；
n
++
）
{
if
（
fabs
（
z
［
n
］
-
c
）
>
EPSILON
）
{
has_error
=
true
；
}
}
printf
（
“%s
\n
”
，
has_error
？
“Has errors”
：
“No errors”
）；
}
編譯後執行得到輸出：
可見宏函式捕抓到了執行時刻錯誤。上面這個意思是指令碼檔案中第30行的cudaMemcpy的引數有問題，因為我們故意將cudaMemcpyHostToDevice寫成了cudaMemcpyDeviceToHost。這樣我們就得到了有用的錯誤資訊。
檢查核函式
核函式不返回任何值，因此上述的檢測功能不適用於核函式。但可以在呼叫核函式後加上如下兩條語句：
CHECK
（
cudaGetLastError
（））；
CHECK
（
cudaDeviceSynchronize
（））；
其中
第一條語句的作用是捕抓第二個語句之前的最後一個錯誤，第二條語句的作用是同步主機於裝置
。之所以要同步主機與裝置，是因為
核函式的呼叫是非同步的，即主機發出調用核函式的命令後會立即執行後面的語句
，不會等待核函式執行完畢。需要注意的是，上述同步函式比較耗時，如果在程式的較內淺層迴圈呼叫的話會嚴重降低程式的效能。
為展示對核函式呼叫的檢查，我們稍改上面的程式碼：
#include
“error。cuh”
#include
#include
const
double
EPSILON
=
1。0e-15
；
const
double
a
=
1。23
；
const
double
b
=
2。34
；
const
double
c
=
3。57
；
void
__global__
add
（
const
double
*
x
，
const
double
*
y
，
double
*
z
，
const
int
N
）；
void
check
（
const
double
*
z
，
const
int
N
）；
int
main
（
void
）
{
const
int
N
=
100000000
；
const
int
M
=
sizeof
（
double
）
*
N
；
double
*
h_x
=
（
double
*
）
malloc
（
M
）；
double
*
h_y
=
（
double
*
）
malloc
（
M
）；
double
*
h_z
=
（
double
*
）
malloc
（
M
）；
for
（
int
n
=
0
；
n
<
N
；
++
n
）
{
h_x
［
n
］
=
a
；
h_y
［
n
］
=
b
；
}
double
*
d_x
，
*
d_y
，
*
d_z
；
CHECK
（
cudaMalloc
（（
void
**
）
&
d_x
，
M
））；
CHECK
（
cudaMalloc
（（
void
**
）
&
d_y
，
M
））；
CHECK
（
cudaMalloc
（（
void
**
）
&
d_z
，
M
））；
CHECK
（
cudaMemcpy
（
d_x
，
h_x
，
M
，
cudaMemcpyHostToDevice
））；
CHECK
（
cudaMemcpy
（
d_y
，
h_y
，
M
，
cudaMemcpyHostToDevice
））；
const
int
block_size
=
1280
；
const
int
grid_size
=
（
N
+
block_size
-
1
）
/
block_size
；
add
<<<
grid_size
，
block_size
>>>
（
d_x
，
d_y
，
d_z
，
N
）；
CHECK
（
cudaGetLastError
（））；
CHECK
（
cudaDeviceSynchronize
（））；
CHECK
（
cudaMemcpy
（
h_z
，
d_z
，
M
，
cudaMemcpyDeviceToHost
））；
check
（
h_z
，
N
）；
free
（
h_x
）；
free
（
h_y
）；
free
（
h_z
）；
CHECK
（
cudaFree
（
d_x
））；
CHECK
（
cudaFree
（
d_y
））；
CHECK
（
cudaFree
（
d_z
））；
return
0
；
}
void
__global__
add
（
const
double
*
x
，
const
double
*
y
，
double
*
z
，
const
int
N
）
{
const
int
n
=
blockDim
。
x
*
blockIdx
。
x
+
threadIdx
。
x
；
if
（
n
<
N
）
{
z
［
n
］
=
x
［
n
］
+
y
［
n
］；
}
}
void
check
（
const
double
*
z
，
const
int
N
）
{
bool
has_error
=
false
；
for
（
int
n
=
0
；
n
<
N
；
n
++
）
{
if
（
fabs
（
z
［
n
］
-
c
）
>
EPSILON
）
{
has_error
=
true
；
}
}
printf
（
“%s
\n
”
，
has_error
？
“Has errors”
：
“No errors”
）；
}
執行緒塊大小的最大值為1024，我們為了實驗錯誤檢測程式，將執行配置中的執行緒塊大小改成1280，然後編譯和執行：
可以看到核函式不能成功地呼叫，並且我們成功捕抓到了核函式的錯誤。它告訴我們第36行出現了核函式地執行配置引數設定無效。如果去掉CHECK（cudaDeviceSynchronize（））其實也能成功捕抓上述錯誤資訊。這是因為後面的資料傳輸CHECK（cudaMemcpy（h_z， d_z， M， cudaMemcpyDeviceToHost））能夠起到一種隱式地同步主機與裝置的作用。一般情況下要獲得精確的錯誤資訊還需要
顯式的同步
。例如呼叫cudaDeviceSynchronize（）函式或者臨時將環境變數CUDA_LAUNCH_BLOCKING的值設為1：
$export
CUDA_LAUNCH_BLOCKING
=
1
這樣核函式的呼叫就是同步的了，即主機呼叫一個核函式後必須等核函式執行完才能向下執行
。
用CUDA-MEMCHECK檢查記憶體錯誤
CUDA有一個名為CUDA-MEMCHECK的工具集，包括
memcheck、racecheck、initcheck、synccheck
四個工具，都可以由可執行檔案cuda-memcheck呼叫：
$cuda
-memcheck
-
-tool
memcheck
［options］ app_name ［options］
$cuda
-memcheck
-
-tool
racecheck
［options］ app_name ［options］
$cuda
-memcheck
-
-tool
initcheck
［options］ app_name ［options］
$cuda
-memcheck
-
-tool
synccheck
［options］ app_name ［options］
其中memcheck工具可以簡化成：
$cuda
-memcheck
［options］ app_name ［options］
上一篇:腸鏡做完後需要注意什麼？
下一篇：燙髮過度該怎麼修復？
