1

我必须将预先存在的“仅主机”反向传播实现移植到 CUDA。我认为算法的性质在这里并不重要，所以我不会对它的工作方式做太多解释。不过，我认为重要的是它使用 3 维数组，其所有三个维度都是动态分配的。我使用带有 CUDA 5.0 的 VS2010。我的设备是2.1。原始主机代码可以在这里下载 → http://files.getwebb.org/view-cre62u4d.html

代码要点：

  1. 使用“pattern.h”中的数据结构将成人数据中的模式加载到内存中。
  2. 分配了几个多维数组
  3. 该算法使用之前分配的数组在模式上运行。

如果您想尝试运行代码，请不要忘记修改 kernel.cu 开头的 PATH 常量。我还建议你使用“2”层、“5”个神经元和“0.00001”的学习率。如您所见，这非常有效。“MSE”正在改善。对于那些不知道这个算法是做什么的人，我们简单地说它学习如何根据模式中存在的 14 个变量来预测目标值。“MSE”减少，意味着算法在每个“epoch”之后犯的错误更少。

我花了很长时间尝试在设备上运行此代码。而我仍然没有成功。最后一次尝试是通过简单地复制初始化数组的代码并将算法运行到一个大内核中来完成的。又失败了。该代码可以在那里下载 → http://files.getwebb.org/view-cre62u4c.html

准确地说，以下是与原始仅主机代码的区别：

  • 算法使用的 f() 和 fder() 成为设备 函数。
  • 参数是硬编码的：2 层，5 个神经元，学习率为 0.00001
  • “w” 数组使用固定值 (0.5) 进行初始化，不再是 rand()
  • a 数据结构在设备内存中分配，数据从host内存中的adult.data加载后发送到设备内存

我认为我做了使代码在内核中运行所需的最少修改。“kernel_check_learningData”内核，显示了有关加载到设备内存中的模式的一些信息，证明了以下代码，将模式从主机发送到设备，确实有效：

Data data;
Data* dev_data;
int* dev_t;
double* dev_x;
...
input_adult(PathFile, &data);
...
cudaMalloc((void**)&dev_data, sizeof(Data));
cudaMalloc((void**)&dev_t, data.N * sizeof(int));
cudaMalloc((void**)&dev_x, data.N * data.n * sizeof(double));
// Filling the device with t and x's data.
cudaMemcpy(dev_t, data.t, data.N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_x, data.x, data.N * data.n * sizeof(double), cudaMemcpyHostToDevice);
// Updating t and x pointers into devices Data structure.
cudaMemcpy(&dev_data->t, &dev_t, sizeof(int*), cudaMemcpyHostToDevice);
cudaMemcpy(&dev_data->x, &dev_x, sizeof(double*), cudaMemcpyHostToDevice);
// Copying N and n.
cudaMemcpy(&dev_data->N, &data.N, sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(&dev_data->n, &data.n, sizeof(int), cudaMemcpyHostToDevice);

当读取“w”数组时，它显然在前向阶段开始时失败。我找不到任何解释。

我看到两种可能性：

  1. 将模式发送到设备内存中的代码存在错误，尽管它似乎可以正常工作，并且在开始前进阶段时会进一步引发错误。
  2. CUDA API 的行为不像它应该的那样！

我拼命寻找我的错误很长一段时间。所以我想知道社区是否可以为我提供一些帮助。

谢谢。

4

1 回答 1

1

这是您的代码中的问题，以及为什么它在 64 位机器模式下工作而不是 32 位机器模式。

在您的反向传播内核中，在正向路径中，您有如下代码序列：

/*
* for layer = 0
*/
for (i = 0; i < N[0]; i++) {    // for all neurons i of layer 0
a[0][i] = x[ data->n * pat + i];    // a[0][i] = input i
}

在32位机器模式下（Win32项目，--machine 32正在传给nvcc），发生write的迭代i=7时失败a[0][7]；此写入超出范围。在这一点上，a[0][7]旨在保存一个double值，但由于某种原因，索引使我们超出了界限。

顺便说一句，您可以通过在构建可执行文件的目录中打开命令提示符并运行以下命令来验证这一点：

cuda-memcheck test_bp

假设test_bp.exe是您的可执行文件的名称。cuda-memcheck 可以方便地识别发生越界写入，甚至识别发生它的源行。

那么为什么这超出了界限呢？让我们先看一下内核代码中a[0][]分配的位置：

a[0] = (double *)malloc( N[0] * sizeof(double *) );
                                              ^ oops!!

a[0][]旨在保存double 数据，但您正在分配指针存储。事实证明，在 64 位机器中，两种类型的存储大小相同，因此最终可以正常工作。但是在 32 位机器中，double 指针是 4 个字节，而double 数据是 8 个字节。因此，在 32 位机器中，当我们以 8 字节的数据步长对这个数组进行索引时，我们最终会跑出数组的末尾。

在内核代码的其他地方，您正在为其他“层”分配存储，a如下所示：

a[layer] = (double *)malloc( N[layer] * sizeof(double) );

哪个是对的。我看到原始的“仅限主机”代码似乎也包含此错误。该代码中也可能存在潜在缺陷。

如果您想在 Windows wddm 设备上运行，您仍然需要以某种方式解决内核运行时间以避免 Windows TDR 事件。正如我已经指出的，这段代码没有尝试使用机器的并行能力。

于 2013-04-30T04:21:57.637 回答