Reply to comment

Как лучше организовать чтение в общую память вот в таком случае:

Tagged:  

Как лучше организовать чтение в общую память вот в таком случае:
Есть матрица t[nx][ny] заведомо больше размера блока(16х16), Каждому потоку требеются
текущее значение t[x][y] и 4 соседа t[x][y+1], t[x][y-1], t[x+1][y], t[x-1][y]. Т.е. в идеале каждый поток читает свое значение в общую память потом делаем синхронизацию блока(те все как в примере с матрицой из sdk). Однако тут получается еще нужны "краевые" значения для блока

Можете что нить лучше предложить такого варианта

  1. __shared__  float sTimecc[BLOCK_SIDE+2][BLOCK_SIDE+2];
  2. const int threadX = threadIdx.x+1;
  3. const int threadY = threadIdx.y+1;
  4.  
  5. sTimecc[threadX][threadY]=t[(y)*nz + (z)];
  6.             if(threadX == blockDim.x){
  7.                 sTimecc[threadX+1][threadY] = t[(y+1)*nz + (z)];
  8.             }
  9.             if(threadX == 1){
  10.                 sTimecc[threadX-1][threadY] = t[(y-1)*nz + (z)];
  11.             }
  12.             if(threadY == blockDim.y){
  13.                 sTimecc[threadX][threadY+1] = t[(y)*nz + (z+1)];
  14.             }
  15.             if(threadY == 1){
  16.                 sTimecc[threadX][threadY-1] = t[(y)*nz + (z-1)];
  17.             }
  18. __syncthreads();

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

CAPTCHA
This question is for testing whether you are a human visitor and to prevent automated spam submissions.
Image CAPTCHA
Enter the characters shown in the image.

Copyright © 2008-2009 Alex Tutubalin