У меня есть шаблонизированное ядро CUDA для расчета и установки значений на границе между двумя вычислительными сетками.Значения рассчитываются с использованием 3-х отдельных вкладов, полученных из функций-членов класса с экземплярами класса, переданными в ядро.Если я получу какой-либо из этих вкладов в одиночку для установки в выводе, ядро будет работать.Как только я добавляю 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;
Я совершенно озадачен, как это кажетсябросить вызов логике и действительно хотел бы понять, что происходит.Кто-нибудь сталкивался с чем-то подобным или с какой-либо идеей, что может быть причиной этого?
Я использую CUDA 9.2 с Visual Studio 2017 и тестировал код на GTX 980 Ti (Compute 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);
}
}
Это похоже на то, как если бы ядра содержали слишком много строк кода (в приведенном выше ядре есть дополнительный кодв различных используемых функциях) не запускается при определенных условиях.