Ядро CUDA, похоже, игнорирует оператор «if»

Question

Ниже приведена часть моего ядра, которое не ведет себя правильно, а затем объяснение того, что я нашел во время отладки.

__global__ void Mangler(float *matrix, int *map) 
{ 
    __shared__ signed int localMap[N]; 

    if(0 == threadIdx.x) 
    { 
     for(int i=0; i<N; i++) 
      localMap[i] = -1; 
    } 

    __syncthreads(); 

    int fn = ...; // a lot of code goes into this number, skipped for clarity 
    int rnumber = threadIdx.x; 

    int X = atomicCAS(&localMap[fn], -1, rnumber); // Spot of bother 1 

    if(X == -1) // Spot of bother 2 
    { 
     // some code 
    } 
    else 
    { 
     // other code 
    } 
} 

Я нашел в документации, atomicCAS(*address, compare, value) в основном возвращает (и сохраняет по данному адресу) результат (old == compare ? value : old), где старое значение по адресу, перед выполнением функции.

Идущий с этим, я считаю, что выполнение int X = atomicCAS(&localMap[fn], -1, rnumber); должны иметь два возможных исхода (согласно Руководству по программированию NVidia Cuda C):

• если localMap[fn] == -1 то X должен иметь значение rnumber и localMap[fn] должен иметь значение от rnumber. Этого не происходит.
• если localMap[fn] != -1 тогда X должно быть установлено в значение localMap[fn], и указанное значение должно быть оставлено без изменений.

Что происходит вместо этого, как и отладка с Nsight показал мне, что X в настоящее время присваивается -1, в то время как в настоящее время localMap[fn] присваивается значение rnumber. Я этого не понимаю, но, как вы можете видеть в моем коде, я изменил if, чтобы поймать эту ситуацию.

подводит меня к пятну хлопоты номера 2: хотя Nsight показывает значение X как -1, if {} не быть полностью пропущен (без контрольных точек в пределах попадания бы то ни было), и выполнение переходит прямо к else.

Мои вопросы:

• ли я неправильно atomicCAS полностью?да, я сделал
• Что может быть причиной и if, которые должны оцениваться как истинные, чтобы прыгать прямо в else в коде устройства?

Я использую NVidia CUDA 5.5, Visual Studio 2012 x64 в Windows 8, NVidia Nsight Monitor Visual Studio Edition 3.1. GPU для машины - NVidia GeForce GTX 550 Ti.

Я попытался изменить синтаксис на if(X!=-1); истинная ветвь if все еще не выполняется.

Answer 1

Из документа atomicCAS возвращается старое значение, то есть в вашем списке ваши два результата неверны. Для вашего X всегда будет установлено старое значение localMap[fn], независимо от того, какое значение оно имело. То, что установлено в соответствии с сравнением с -1, является новым значением localMap[fn]. Когда оно равно -1, оно установлено в rnumber, иначе оно остается нетронутым.

Таким образом, поведение, которое вы видите со значениями X, rnumber и localMap, как ожидается.

Я не могу помочь с вашей второй проблемой, так как я не использую NSight и не знаю, как это работает - в соответствии с вашим кодом ваша истинная ветка должна быть оценена (но будьте осторожны: ваша ложная ветка также - threaded some threads может иметь условие, которое оценивается как true, а some - false - мое предположение/предположение состояло в том, что вы должны каким-то образом указать ваш отладчик, какой поток/warp/block вы хотите отлаживать, и вы посмотрели на false).