Проблемы при работе с текстурами

Помогите, кто чем может!
Вобщем есть 2 алгоритма. Первый берет на вход двумерный массив, и заполняет другой двумерный массив. Второй берет на вход результат работы предыдущего (2й массив) и заполняет 3й.
Все массивы 2мерные, представлены как одномерные в глобальной памяти. Все работает.
Когда решил работать с ними, как с текстурами, возникли проблемы...

Создание 1го массива:

  1. CUDA_ARRAY_DESCRIPTOR desc;
  2. desc.Format      = CU_AD_FORMAT_SIGNED_INT32;
  3. desc.NumChannels = 1;
  4. desc.Width       = image.Width;    //  1280
  5. desc.Height      = image.Height;  //  750
  6. cuArrayCreate(&arrImage,&desc);
  7. ...
  8. cuTexRefSetArray(texImage,arrImage,CU_TRSA_OVERRIDE_FORMAT);
  9. cuTexRefSetAddressMode(texImage,0,CU_TR_ADDRESS_MODE_WRAP);
  10. cuTexRefSetAddressMode(texImage,1,CU_TR_ADDRESS_MODE_WRAP);
  11. cuTexRefSetFilterMode(texImage,CU_TR_FILTER_MODE_LINEAR);
  12. cuTexRefSetFlags(texImage,CU_TRSF_READ_AS_INTEGER);
  13. cuTexRefSetFormat(texImage,CU_AD_FORMAT_SIGNED_INT32,1);

Запускаю 1й алгоритм все прекрасно работает.
Второй массив создаю из уже существующего в глобальной памяти:

  1. cuMemAlloc(&d_Rastr,rastr.Height*rastr.Width*sizeof(Rastr::SMEl));
  2. ...
  3. CUDA_ARRAY_DESCRIPTOR desc;
  4. desc.Format      = CU_AD_FORMAT_SIGNED_INT32;
  5. desc.NumChannels = 1;
  6. desc.Width       = rastr.Width;    //   1280
  7. desc.Height      = rastr.Height;  //   749
  8. cuTexRefSetAddress2D(texRastr,&desc,d_Rastr,rastr.Width*sizeof(Rastr::SMEl));
  9. cuTexRefSetAddressMode(texRastr,0,CU_TR_ADDRESS_MODE_WRAP);
  10. cuTexRefSetAddressMode(texRastr,1,CU_TR_ADDRESS_MODE_WRAP);
  11. cuTexRefSetFilterMode(texRastr,CU_TR_FILTER_MODE_LINEAR);
  12. cuTexRefSetFlags(texRastr,CU_TRSF_READ_AS_INTEGER);

Запускаю 2й алгоритм - все виснет, экран гаснет... В чем проблема может быть??

Вот код 2го алгоритма:

  1. texture<int, 2, cudaReadModeElementType> Rastr;
  2. ...
  3. extern "C" __global__ void GetRects
  4. (int           rWidth,
  5.  int           rHeight,
  6.  int2         *rects,
  7.  unsigned int *len)
  8. {
  9.  int x = blockIdx.x*blockDim.x+threadIdx.x;
  10.  int y = blockIdx.y*blockDim.y+threadIdx.y;
  11.  
  12.  if(x>rWidth || y>rHeight)
  13.    return;
  14.  
  15.  if(tex2D(Rastr,x+0.5f,y+0.5f) == 0)
  16.    return;
  17.  
  18.  unsigned int n;
  19.  int lXY,_x,_y = y,pre = 32767,val;
  20.  // Идем по строкам
  21.  SetXY(lXY,x,y);
  22.  do
  23.    {
  24.         // Идем по столбцам
  25.         _x = x;
  26.         do
  27.          {
  28.           _x++;
  29.           val = tex2D(Rastr,x+0.5f,y+0.5f);
  30.          }
  31.         while(_x<pre && (val & 5)==0);
  32.  
  33.         // Запоминаем прямоугольник
  34.         if(_y!=y && pre>_x)
  35.           {
  36.            n = atomicInc(len,MAX_INT);
  37.            rects[n].x = lXY;
  38.            SetXY(rects[n].y,pre,y);
  39.           }
  40.         pre = _x;
  41.         y++;
  42.         val = tex2D(Rastr,x+0.5f,y+0.5f);
  43.    }
  44.  while((val & 10) == 0);
  45.  // Запоминаем прямоугольник
  46.  if(pre == _x)
  47.    {
  48.         n = atomicInc(len,MAX_INT);
  49.         rects[n].x = lXY;
  50.         SetXY(rects[n].y,_x,y);
  51.    }
  52. }

Forums: 

Секундная задержка где-то,

Секундная задержка где-то, потом темный экран, еще через секунду винда драйвер восстанавливает.

Скорее всего дело не в этом,

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

Вопрос в том, сходится ли

Вопрос в том, сходится ли он.

У вас признак завершения - наличие битика 0x10 в val, фетчится ли все из текстуры правильно (проверка понятная - переписать из текстуры в global и посмотреть, получено ли ожидаемое).

Поставил серию экспериментов,

Поставил серию экспериментов, результаты печальные. Текстуры работают раз в 5 медленнее чем глобальная память. В этом и было дело видимо.

Вот код работающей

Вот код работающей ф-и.

  1. texture<int, 2, cudaReadModeElementType> Image;
  2. ...
  3. extern "C" __global__ void CreateRastr(int  iWidth,
  4.                                        int  iHeight,
  5.                                        int *rastr,
  6.                                        int  rWidth)
  7. {
  8.  int x = blockIdx.x*blockDim.x+threadIdx.x;
  9.  int y = blockIdx.y*blockDim.y+threadIdx.y;
  10.  
  11.  if(x>iWidth || y>iHeight)
  12.    return;
  13.    
  14.  int c11,c12,c21,c22,b=0;
  15.  
  16.  c11 = tex2D(Image,x+0.5f,y+0.5f);
  17.  c12 = tex2D(Image,x+1+0.5f,y+0.5f);
  18.  c21 = tex2D(Image,x+0.5f,y+1+0.5f);
  19.  c22 = tex2D(Image,x+1+0.5f,y+1+0.5f);
  20.  
  21.  if(c12 != c22) b|=8; // Вправо
  22.  if(c12 != c11) b|=4; // Вверх
  23.  if(c21 != c11) b|=2; // Влево
  24.  if(c21 != c22) b|=1; // Вниз
  25.  
  26.  rastr[y*rWidth+x] = b;
  27. }
  28. //-------------------------------------------------------------------------------------------