在内核中进行简单的代码更改后，CUDA 内核无法启动

Question

我有一个模板化的 CUDA 内核，用于在 2 个计算网格之间的接口处计算和设置值。这些值是使用 3 个单独的贡献计算的，这些贡献是从 class 成员函数中获得的，其中 class 个实例传递给内核。如果我单独获得这些贡献中的任何一个以在输出中设置内核工作。一旦我将这些贡献中的 2 个（或全部）添加到输出中，内核就根本不会启动。

我在最后插入了完整的内核代码，但我会先尝试举例说明上面的代码。

首先定义前2个贡献：

//contribution 1
VType value1 = (V_m2 * 2 * b_val_sec / 3 + V_2 * (b_val_pri + b_val_sec / 3)) / (b_val_sec + b_val_pri);
//contribution 2
VType value2 = (Vdiff2_sec * b_val_sec * hL * hL - Vdiff2_pri * b_val_pri * hR * hR) / (b_val_sec + b_val_pri);

现在设置输出：

案例 1 - 内核启动并设置预期值：

V_pri[cell1_idx] = value1;

案例 2 - 内核启动并设置预期值：

V_pri[cell1_idx] = value2;

情况 3 - 内核未启动：

V_pri[cell1_idx] = value1 + value2;

我完全被难住了，因为这似乎违背了逻辑，我真的很想了解正在发生的事情。有没有人遇到过类似的事情，或者知道是什么原因造成的？

我在 Visual Studio 2017 中使用 CUDA 9.2，我在 GTX 980 Ti（计算 5.2）和 GTX 1060（计算 6.1）上测试了代码，结果相同。

这里是完整的内核代码：

template <typename VType, typename Class_CMBND>
__global__ void set_cmbnd_values_kernel(
cuVEC_VC<VType>& V_sec, cuVEC_VC<VType>& V_pri,
Class_CMBND& cmbndFuncs_sec, Class_CMBND& cmbndFuncs_pri,
CMBNDInfoCUDA& contact)
{

 int box_idx = blockIdx.x * blockDim.x + threadIdx.x;

 cuINT3 box_sizes = contact.cells_box.size();

 if (box_idx < box_sizes.dim()) {

     int i = (box_idx % box_sizes.x) + contact.cells_box.s.i;
     int j = ((box_idx / box_sizes.x) % box_sizes.y) + contact.cells_box.s.j;
     int k = (box_idx / (box_sizes.x * box_sizes.y)) + contact.cells_box.s.k;

     cuReal hL = contact.hshift_secondary.norm();
     cuReal hR = contact.hshift_primary.norm();
     cuReal hmax = (hL > hR ? hL : hR);

     int cell1_idx = i + j * V_pri.n.x + k * V_pri.n.x*V_pri.n.y;

     if (V_pri.is_empty(cell1_idx) || V_pri.is_not_cmbnd(cell1_idx)) return;

     int cell2_idx = (i + contact.cell_shift.i) + (j + contact.cell_shift.j) * V_pri.n.x + (k + contact.cell_shift.k) * V_pri.n.x*V_pri.n.y;

     cuReal3 relpos_m1 = V_pri.rect.s - V_sec.rect.s + ((cuReal3(i, j, k) + cuReal3(0.5)) & V_pri.h) + (contact.hshift_primary + contact.hshift_secondary) / 2;

     cuReal3 stencil = V_pri.h - cu_mod(contact.hshift_primary) + cu_mod(contact.hshift_secondary);

     VType V_2 = V_pri[cell2_idx];
     VType V_m2 = V_sec.weighted_average(relpos_m1 + contact.hshift_secondary, stencil);

     //a values
     VType a_val_sec = cmbndFuncs_sec.a_func_sec(relpos_m1, contact.hshift_secondary, stencil);
     VType a_val_pri = cmbndFuncs_pri.a_func_pri(cell1_idx, cell2_idx, contact.hshift_secondary);

     //b values adjusted with weights
     cuReal b_val_sec = cmbndFuncs_sec.b_func_sec(relpos_m1, contact.hshift_secondary, stencil) * contact.weights.i;
     cuReal b_val_pri = cmbndFuncs_pri.b_func_pri(cell1_idx, cell2_idx) * contact.weights.j;

     //V'' values at cell positions -1 and 1
     VType Vdiff2_sec = cmbndFuncs_sec.diff2_sec(relpos_m1, stencil);
     VType Vdiff2_pri = cmbndFuncs_pri.diff2_pri(cell1_idx);

     //Formula for V1
     V_pri[cell1_idx] = (V_m2 * 2 * b_val_sec / 3 + V_2 * (b_val_pri + b_val_sec / 3)
        - Vdiff2_sec * b_val_sec * hL * hL - Vdiff2_pri * b_val_pri * hR * hR
        + (a_val_pri - a_val_sec) * hmax) / (b_val_sec + b_val_pri);
 }
}

在某些情况下，内核代码行数过多（在上述内核中使用的各种函数中有额外的代码）几乎无法启动。

Answer 1

好的，看来我找到了问题的答案。

查看生成的错误，我得到 "Too many resources requested for launch"。

我已经将每个块的线程数从 512 减少到 256，内核现在运行良好。

在内核中进行简单的代码更改后，CUDA 内核无法启动

CUDA kernel fails to launch after making simple code changes inside the kernel

cuda

scientific-computing