CUDA nvcc неправильно компилирует тернарные операторы/условные сокращения?

РЕДАКТИРОВАНИЕ
Вот небольшая программа, которую вы скомпилируете, чтобы увидеть такие ошибки самостоятельно...

//for printf
#include <stdio.h>

#include <cuda.h>

__inline __host__ void gpuAssert(cudaError_t code, char *file, int line, 
                 bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code),
          file, line);
      //if (abort) exit(code);
   }
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

__global__ void myKernel1(int *dev_idx, int *dev_tID, const int offset)
{
   int myElement = threadIdx.x + blockDim.x * blockIdx.x;
   //
   int temp;
   temp = myElement+
      offset +
      (offset==0)?0:(offset&0x01==0x0)?(offset-1)*(offset>>1):
      (offset)*(offset>>1);
   dev_idx[myElement+offset] = temp;
   dev_tID[myElement+offset] = myElement;

}

__global__ void myKernel2(int *dev_idx, int *dev_tID, const int offset)
{
   int myElement = threadIdx.x + blockDim.x * blockIdx.x;
   //
   int temp;
   temp = myElement+offset;
   if (offset != 0 && offset&0x01==0x0) temp+= (offset-1)*(offset>>1);
   if (offset != 0 && offset&0x01!=0x0) temp+= offset*( offset>>1);
   dev_idx[myElement+offset] = temp;
   dev_tID[myElement+offset] = myElement;

}

__host__ void PrintMethod1(int *h_idx, int * h_tID, const int offset, 
               const int total, const int h_set)
{
   for (int c=(h_set==0)?0:offset;
    c < (h_set==0)?offset:total;
    c++)
      printf("Element #%d --> idx: %d   tID: %d\n",
         c,h_idx[c],h_tID[c]);
}

__host__ void PrintMethod2(int *h_idx, int * h_tID, const int offset, 
               const int total, const int h_set)
{
   int loopStart = (h_set==0)?0:offset;
   int loopEnd = (h_set==0)?offset:total;
   printf("Loop Start: %d, Loop End: %d\n",
      loopStart, loopEnd);
   for (int c=loopStart; c < loopEnd; c++)
      printf("Element #%d --> idx: %d   tID: %d\n",
         c,h_idx[c],h_tID[c]);
}

//Checks if there is a compatible device
bool IsCompatibleDeviceRunning()
{
   int *dummy;
   return cudaGetDeviceCount(dummy) != cudaSuccess;
}

int main()
{
   //Check for compatible device
   if (!IsCompatibleDeviceRunning())
   {
      printf("ERROR: No compatible CUDA devices found!\n");
      exit(1);
   }
   const int total = 30;
   const int offset = total/2;

   int * h_tID, * dev_tID, * h_idx, * dev_idx, h_set;
   h_tID = (int *) malloc(total*sizeof(int));
   h_idx = (int *) malloc(total*sizeof(int));
   gpuErrchk(cudaMalloc((void **) &dev_tID,total*sizeof(int)));
   gpuErrchk(cudaMalloc((void **) &dev_idx,total*sizeof(int)));
   myKernel1<<<1,offset>>>(dev_idx, dev_tID, 0);
   //myKernel2<<<1,offset>>>(dev_idx, dev_tID, 0);
   gpuErrchk(cudaPeekAtLastError());
   gpuErrchk(cudaDeviceSynchronize());
   myKernel1<<<1,offset>>>(dev_idx, dev_tID, offset);
   //myKernel2<<<1,offset>>>(dev_idx, dev_tID, offset);
   gpuErrchk(cudaPeekAtLastError());
   gpuErrchk(cudaDeviceSynchronize());
   gpuErrchk(cudaMemcpy(h_tID, dev_tID, total*sizeof(int),
            cudaMemcpyDeviceToHost));
   gpuErrchk(cudaMemcpy(h_idx, dev_idx, total*sizeof(int),
            cudaMemcpyDeviceToHost));
   h_set = 0;
   //PrintMethod1(h_idx, h_tID, offset, total, h_set);
   PrintMethod2(h_idx, h_tID, offset, total, h_set);
   h_set = 1;
   //PrintMethod1(h_idx, h_tID, offset, total, h_set);
   PrintMethod2(h_idx, h_tID, offset, total, h_set);
   return 0;
}

При запуске MyKernel2 в массив записывается правильный вывод:

Loop Start: 0, Loop End: 15
Element #0 --> idx: 0   tID: 0
Element #1 --> idx: 1   tID: 1
Element #2 --> idx: 2   tID: 2
Element #3 --> idx: 3   tID: 3
Element #4 --> idx: 4   tID: 4
Element #5 --> idx: 5   tID: 5
Element #6 --> idx: 6   tID: 6
Element #7 --> idx: 7   tID: 7
Element #8 --> idx: 8   tID: 8
Element #9 --> idx: 9   tID: 9
Element #10 --> idx: 10   tID: 10
Element #11 --> idx: 11   tID: 11
Element #12 --> idx: 12   tID: 12
Element #13 --> idx: 13   tID: 13
Element #14 --> idx: 14   tID: 14
Loop Start: 15, Loop End: 30
Element #15 --> idx: 120   tID: 0
Element #16 --> idx: 121   tID: 1
Element #17 --> idx: 122   tID: 2
Element #18 --> idx: 123   tID: 3
Element #19 --> idx: 124   tID: 4
Element #20 --> idx: 125   tID: 5
Element #21 --> idx: 126   tID: 6
Element #22 --> idx: 127   tID: 7
Element #23 --> idx: 128   tID: 8
Element #24 --> idx: 129   tID: 9
Element #25 --> idx: 130   tID: 10
Element #26 --> idx: 131   tID: 11
Element #27 --> idx: 132   tID: 12
Element #28 --> idx: 133   tID: 13
Element #29 --> idx: 134   tID: 14

Когда MyKernel1 запускается с идентичным присвоением idx на основе троичной системы, он получает ноль для всех результатов:

Loop Start: 0, Loop End: 15
Element #0 --> idx: 0   tID: 0
Element #1 --> idx: 0   tID: 1
Element #2 --> idx: 0   tID: 2
Element #3 --> idx: 0   tID: 3
Element #4 --> idx: 0   tID: 4
Element #5 --> idx: 0   tID: 5
Element #6 --> idx: 0   tID: 6
Element #7 --> idx: 0   tID: 7
Element #8 --> idx: 0   tID: 8
Element #9 --> idx: 0   tID: 9
Element #10 --> idx: 0   tID: 10
Element #11 --> idx: 0   tID: 11
Element #12 --> idx: 0   tID: 12
Element #13 --> idx: 0   tID: 13
Element #14 --> idx: 0   tID: 14
Loop Start: 15, Loop End: 30
Element #15 --> idx: 0   tID: 0
Element #16 --> idx: 0   tID: 1
Element #17 --> idx: 0   tID: 2
Element #18 --> idx: 0   tID: 3
Element #19 --> idx: 0   tID: 4
Element #20 --> idx: 0   tID: 5
Element #21 --> idx: 0   tID: 6
Element #22 --> idx: 0   tID: 7
Element #23 --> idx: 0   tID: 8
Element #24 --> idx: 0   tID: 9
Element #25 --> idx: 0   tID: 10
Element #26 --> idx: 0   tID: 11
Element #27 --> idx: 0   tID: 12
Element #28 --> idx: 0   tID: 13
Element #29 --> idx: 0   tID: 14

Когда запускается PrintMethod1 (с тернарным ограничением), он segfaults, по существу застревая в бесконечном цикле. Обратите внимание, это на стороне хоста!!

При запуске PrintMethod2 выходные данные обычно выводятся, как и ожидалось выше.

Вот моя команда компиляции:

nvcc --compiler-options -fno-strict-aliasing -DUNIX -m64 -O2 \
--compiler-bindir /usr/bin/g++ \
-gencode=arch=compute_20,code=\"sm_21,compute_20\" \
-I/usr/local/CUDA_SDK/C/common/inc -I/usr/local/CUDA_SDK/shared/inc \
-o TEST Test.cu

Единственная подсказка, которая у меня есть, заключается в том, что он жалуется на то, что оба ядра имеют неправильный параметр, хотя он выглядит правильно и дает правильные результаты для MyKernel2.

Я думаю, что приведенный выше пример — это в значительной степени то, что комментаторы могли бы попробовать самостоятельно, основываясь на приведенном ниже описании, но он экономит ваше время и усилия на написание кода!

Дайте мне знать, если есть что-то еще, что я могу опубликовать, чтобы помочь разобраться.

Исходный вопрос

Большинство компиляторов C, как определено яз. стандартная поддержка тернарных операторов.

e.g.

int myVar;
myVar=(testFlg==true)?-1:1;

Однако, на удивление, CUDA nvcc, по-видимому, удаляет некоторые тернарные операторы и заменяет их нулями, когда они используются в ядре...

Я обнаружил это, применив cuPrintf для проверки проблемного блока кода. Например, предположим, что у меня есть два ядра, совместно использующие глобальный массив для вывода. Первое ядро ​​имеет дело с первой порцией элементов. Второе ядро ​​получает смещение, указывающее, как далеко нужно перейти в массиве, чтобы не перезаписать элементы первого ядра. Смещение отличается для четных и нечетных.

Так что я мог бы написать:

if (krnl!=0 && offset&0x01==0x0)
   idx+=(offset-1)*(offset>>1);
if (krnl!=0 && offset&0x01!=0x0)
   idx+=offset*(offset>>1);

Но было бы более компактно и читабельно (на мой взгляд) написать почти эквивалентный сокращенный синтаксис.

idx += (krnl==0)?0:(offset&0x01==0)?
   (offset-1)*(offset>>1):
   offset*(offset>>1);

Последний код, тем не менее, всегда будет давать ноль, так как компилятор CUDA отсекает сокращенные условные операторы.

Я понимаю, что этим кодом функции злоупотребляют и вызывают расхождение потоков, но в простых случаях кажется, что он не будет отличаться от стандартных условных выражений, если компилятор обработает его правильно.

Это ошибка компилятора или он намеренно не поддерживается?

Кто-нибудь знает, появится ли эта функция в CUDA?

Я был очень удивлен, узнав, что это было источником моих сбоев и ошибок в адресации...

EDIT
Это стандартная функция C, я неправильно понял и ошибочно сказал, что это нестандартно.

РЕДАКТИРОВАТЬ 2
Я сказал "задыхается и умирает" для компилятора. Термин «умирает» определенно неуместен. Скорее, nvcc завершает компиляцию, но, по-видимому, убрал тернарное присваивание на основе оператора и заменил его нулем. Позже это вернется и укусит меня, поскольку данные не записывались в нужные места, а эти места, в свою очередь, использовались в качестве индексов в схеме двойной индексации. Индексы использовались во время завершения на стороне ЦП, поэтому ошибка сегментации произошла на стороне ЦП, но была вызвана фрагментацией компилятора.

Я использую компилятор версии 4.1 и -O2 включен. Похоже, что оптимизатор может оптимизировать переменную, используемую внутри троичной операции, которая может быть источником этой ошибки.

Подверженная ошибкам троичная операция почти идентична примеру I. дал выше, но участвует в большой операции добавления.

Я планирую последовать совету комментатора ниже и отправить отчет об ошибке в NVIDIA, но я оставляю этот пост в качестве предупреждения для других.

Изменить 3

Вот слегка очищенный полный оператор, который всегда дает ноль:

__global__ void MyFunc
( const int offset
  const CustomType * dev_P,
  ...
  const int box)
{
   int tidx = blockIdx.x * blockDim.x + threadIdx.x;
   int idx=0;
   ...
   idx = tidx +
      dev_P->B +
      (box == 0)?0:(offset&0x01!=0x0):
      (offset-1)*(offset>>1):offset*(offset>>1);
   //NOTES:
   //I put the cuPrintf here.... from it I could see that tidx was diff. ints (as you 
   //would expect), but that when added together the sum was always "magically"
   //becoming zero.  The culprit was the nested ternary operator.
   //Once I replaced it with the equivalent conditional, the assignment worked as
   //expected.
   //"offset" is constant on the level of this kernel, but it is not always 0.
   //Outside the kernel "offset" varies greatly over the course of the simulation,
   //meaning that each time the kernel is called, it likely has a different value.
   //"tidx" obviously varies.
   //but somehow the above sum gave 0, likely due to an unreported compiler bug.
   //box is either 0 or 1.  For a certain type of op in my simulation I call this 
   //kernel twice, once for box value 0 and a second time for box value 1
   ...
}