Home :: IDE для CUDA

Reply to comment

Пока что хочу разобраться без
Submitted by blackswan on Fri, 11/13/2009 - 13:00.

Пока что хочу разобраться без оптимизации. Оказалось, что функция работает, однако только для маленьких изображений, т.е. проверял на нескольких, в итоге на изображении в 400Х400 и выше уже не работает, совсем не запускается, а на тех, что меньше, запускается и туда заходит. Как это связано с размером, не могу понять. Вот еще раз код функции

  1. __global__ void __prod(float *A, float *B, long wA, long wB, float *C)
  2. {
  3.         long bx = blockIdx.x;
  4.         long by = blockIdx.y;
  5.         long tx = threadIdx.x;
  6.         long ty = threadIdx.y;
  7.         long aBegin = wA * BLOCK_SIZE * by;
  8.         long aEnd = aBegin + wA - 1;
  9.         long aStep = BLOCK_SIZE;
  10.         long bBegin = BLOCK_SIZE * bx;
  11.         long bStep = BLOCK_SIZE * wB;
  12.         float Csub = 0;
  13.         for (long a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep)
  14.         {
  15.         __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
  16.         __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
  17.         As[ty][tx] = A[a + wA * ty + tx];
  18.         Bs[ty][tx] = B[b + wB * ty + tx];
  19.         __syncthreads();
  20.         for (long k = 0; k < BLOCK_SIZE; ++k) Csub += As[ty][k] * Bs[k][tx];
  21.         __syncthreads();
  22.         }
  23.         long c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;
  24.         C[c + wB * ty + tx] = Csub;
  25. }
  26. void prod(float *&d_A, float *&d_B, float *&d_C, long wA, long hA, long wB)
  27. {
  28.         dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
  29.         dim3 dimGrid(wB / dimBlock.x, hA / dimBlock.y);
  30.         __prod<<<dimGrid, dimBlock>>>(d_A, d_B, wA, wB, d_C);
  31. }
  32.  
  33. вот так я делаю преобразование
  34. void rgb_yiq(float *&img_in, float *&img_out, long iH,long iW, bool dir = true )
  35. {
  36.     if (dir)
  37.     {
  38.         float RGBYIQH[3][3]={{0.299 , 0.587 ,0.114},{ 0.595716, -0.274453, -0.321263},{ 0.211456, -0.522591, 0.311135}};
  39.         float *RGBYIQ;
  40.         cudaMalloc((void**) &RGBYIQ,3*3*sizeof(float));
  41.         cudaMemcpy(RGBYIQ,RGBYIQH,3*3*sizeof(float),cudaMemcpyHostToDevice);
  42.         prod(RGBYIQ,img_in,img_out,3,3,iH*iW);
  43.         cudaFree(RGBYIQ);
  44.     }
  45.     else
  46.     {
  47.         float YIQRGBH[3][3] = {{1.0, 0.9563, 0.621},{ 1.0, -0.2721, -0.6474},{ 1.0, -1.107, 1.7046}};
  48.         float *YIQRGB;
  49.         cudaMalloc((void**) &YIQRGB,3*3*sizeof(float));
  50.         cudaMemcpy(YIQRGB,YIQRGBH,3*3*sizeof(float),cudaMemcpyHostToDevice);
  51.         prod(YIQRGB,img_in,img_out,3,3,iH*iW);
  52.         cudaFree(YIQRGB);
  53.     }
  54. }

Reply

The content of this field is kept private and will not be shown publicly.
  • Web page addresses and e-mail addresses turn into links automatically.
  • Allowed HTML tags: <a> <em> <strong> <cite> <code> <ul> <ol> <li> <dl> <dt> <dd> <i> <table> <td> <tr> <th>
  • Lines and paragraphs break automatically.
  • You can enable syntax highlighting of source code with the following tags: <code>, <blockcode>. The supported tag styles are: <foo>, [foo].
  • Images can be added to this post.

More information about formatting options

[x]
Искать: 

Вход

Connect
Sign in using Facebook
Copyright © 2008-2011 Alex Tutubalin