见到了磅礴的妖魔海，勇者一定会被本魔王强大的力量震慑吧！

不过，当务之急，是找个安全的地方修炼。要是魔功走火入魔，威力大减，那可就大大的不好。

好在我与黑魔法之渊库达，呼吸与共，灵力相通。几番试探，我已了解黑魔法之渊种种禁忌。若有违反，轻则召唤失败，重则反噬魔王本营，再凶猛的魔王，也得谨慎行事！

$\Large\color{blue}{【召唤三诫】}$

✨ $\color{Gold}{有限魔物}：同时活动的魔物数量有限!$

✨ $\color{Gold}{契约访问}：不可访问请求之外的土地$!％

✨ $\color{Gold}{愚昧实体}：关于本大地的一切继承与依附在cuda中均不可用，它们也不支持可变的内存，在早期不能互相组合。$

三诫从魔王无数次尝试得来，要一一理解还为时过早。所有预言和箴言，在后续章节才会一一解释。

$\Large\color{pink}{【谦卑信使】}$

cuda的所有请求都不会被咆哮着拒绝。它会静默地放弃所有的不合理请求，就像渗入勇者团队的史莱姆和哥布林。

所以，作为魔王，首先要怀以谦卑的内心，为每次请求解析cuda的谜语，了解真实情况到底如何。

#define CHECK(call)

do \

{ \

const cudaError_t error_code = call;\

if (error_code != cudaSucess) \

{ \

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）

这样，每次向cuda请法，本魔王都能得知是否触犯戒律。

譬如, 本王要把一切纳入麾下：

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

#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)

int main(void) {
    const int N = 9600000;
    const int M = 9600000;
    double *d_x;
    CHECK(cudaMalloc((void**)&d_x, 1000000000000)); //!!!!!!
    CHECK(cudaFree(d_x));
    CHECK(cudaDeviceSynchronize());
}

却遭到了冷酷无情的否定:

CUDA Error: 
 File: c.cu
 Line: 23
 Error code: 2
 Error text: out of memory

没有足够的空间容纳本王的万万亿计划，可惜！

$\Large\color{gray}{【无主之魂】}$

在cuda眼中，我的小怪不过是些无关紧要的【核函数】，像路边的顽石，滚到哪算哪，就算密语也无法解出。

此时，便需要再笼统地询问一次：

CHECK(cudaGetLastError());

CHECK(cudaDeviceSynchronize());

如果它们出事了，求求告诉本王，本王必定严加管教。

譬如我曾经沉迷回文魔法，要121 * 12321 个小怪随便计算点什么，测测它们的反应速度。

void global count() {
const int n = threadIdx.x;
auto t1 = clock();
double sum = 0;
for (int i=0; i<10000; i++) {
sum += i;
}
auto t2 = clock();
if (n % 1000 == 0) {
printf(“time cost %llu\n”, t2 - t1);
}
}

int main(void) {
count<<< 121, 12321 >>>();
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
}

结果又遭到了拒绝。

CUDA Error:
File:
Line:
Error code: 9
Error text: invalid configuration argument
```

几经尝试，我才发现12321一栏最大值只能有1024。所以我将两个参数颠倒了下。哈，看来是一些障眼法。不过，所有的障眼法都暗中标好了伤害值。这一点，本王尚未领悟。

$\Large\color{red}{【NPCの正义】}$

不过，经过本王四处游历，cuda大陆有着活跃而正义的NPC团体。它们隶属于CUDA-MEMCHECK工会，共有四名:

• memcheck

• racecheck

• initcheck

• synccheck

各位NPC各司其职。

memcheck: 检查是否访问了未申请的内存范围。

racecheck: 检查多个任务是否争抢同一个共享资源

initcheck: 检测全局资源是否初始化

synccheck: 检测多个任务是否能执行同步操作

犹如天书。

本王只需要memcheck帮忙，申请到魔法，放出属于自己的小怪就好了。