写shared memory比读取要耗费很多时间吗

最近我想做一个自己的cuda版的FFT,先测试了一下cufft,512*512的变换。
数据传输用了大约7,500,000 ticks, 执行变换大约17,000,000 ticks

然后测试了一下自己的代码,涉及到2次全局内存读取,约40次shared memory的读取和写入。发现执行kernel大约
用了11,000,000 ticks.在很神奇库的实现的同时,我也为自己的代码感到不解。

是不是写入要比读取共享内存耗费百倍的时间呢?

#define MATRIX_SIZE (1 < <9) 
#define BLOCK_SIZE (1 < <9) 
__global__ 
void FFT_2D_Radix2(DATA_TYPE* dg_buffer, int N ) 
{ 
int tid, rev, pos, pre, stride = 33; 
tid = threadIdx.x; 
rev = bit_reverse3(tid, tail_zero_nums(N)); 
__shared__ DATA_TYPE s_DataR[MATRIX_SIZE]; // 512*4 = 2kB 
__shared__ DATA_TYPE s_DataI[MATRIX_SIZE]; // 512*4 = 2kB 
__shared__ DATA_TYPE s_CosTable[MATRIX_SIZE]; // 512*4 = 2kB 
__shared__ DATA_TYPE s_SinTable[MATRIX_SIZE]; // 512*4 = 2kB 

pos = tid * stride % MATRIX_SIZE; 
s_DataR[pos] = dg_buffer[blockIdx.x * BLOCK_SIZE + rev];            //------------------------------------time-consuming 
s_DataI[pos] = dg_buffer[N*N + blockIdx.x * BLOCK_SIZE + rev];    //-----------------------------------time-consuming 
float theta = GV_2PI / N; 
s_SinTable[pos] = __sinf( theta * tid ); 
s_CosTable[pos] = __cosf( theta * tid ); 
__syncthreads(); 
int step, w; 
for(step = 1;step <N;step=step*2) 
{ 
  if(tid & step) 
  { 
  w = ( tid & ( step - 1 ) ) * stride % MATRIX_SIZE; 
  DATA_TYPE tempR = s_DataR[pos] * s_CosTable[w] + s_DataI[pos] * s_SinTable[w]; 
  DATA_TYPE tempI = s_DataI[pos] * s_CosTable[w] - s_DataR[pos] * s_SinTable[w]; 
  pre = ( tid - step ) * stride % MATRIX_SIZE; 
  
  s_DataR[pos] = s_DataR[pre] - tempR;    //-----------------------------------------------time-consuming 
  s_DataI[pos] = s_DataI[pre] - tempI;        //-----------------------------------------------time-consuming 
  s_DataR[pre] += tempR;                        //-----------------------------------------------time-consuming 
  s_DataI[pre] += tempI;                          //-----------------------------------------------time-consuming 
  } 
  __syncthreads(); 
} 
}

写入和读取shared memory带宽和延迟都是一样的

这样统计时间是不是存在问题,读数据s_DataR[pos] = dg_buffer[blockIdx.x * BLOCK_SIZE + rev]; 是直接读,写时 s_DataR[pos] = s_DataR[pre] - tempR 却是先读,再计算再写

是啊!
理论上读应该只有读,写只能写,不能包括计算

[

恩!.

External Media

开源图形处理器体系结构论坛(OpenGPU论坛) http://www.opengpu.org/bbs/

OpenGPU Graphics Open Source community图形开源社区),聚焦领域(focus domain)包括:
* GPU Architecture(图形处理器体系结构).
* Graphics Algorithm(图形算法).
* GPGPU Programming (面向通用的图形处理器编程).
* Open Source Rendering Engine(开源渲染器).
* Open Source GPU Simulator/RTL Implement(开源GPU模拟器).

[ 本帖最后由 OpenGPU 于 2010-3-10 17:12 编辑 ]

听大牛说的,没错的:)