0

代码片段计算 3D 矢量幅度： mag= sqrt(X*X + Y*Y + Z*Z)

// Note: blockDim.x = 300, gMem= ptr to global mem chunk
__shared__ sMem[100];
float regA;
for (j=0; j<50; j++) {
  if(threadIdx.x < 3) {
    regA= gMem[j];
    sMem[threadIdx.x]= regA*regA;  // Line A   write 5193
  }  
  __syncthreads();
  if(threadIdx.x == 0) {
    regA= sMem[0];                 // Line B   read 5197
    regA+= sMem[1];
    regA+= sMem[2];                // Line C   read 5199 
    sMem[0]= sqrt(regA);
  }
}

同步线程正在防止 RAW 危害。没有它，我会收到 RAW 和 WAR 危险警告。W /它，我仍然收到WAR警告。警告来自线对 AB 和 AC。

 WARN:(Warp Level Programming) Potential WAR hazard detected at __shared__ 0x30b in block (0, 0, 0) :
     Read Thread (0, 0, 0) at 0x000000b0 in /src/trap.cu:5199:Mag(float const *, float const *, int, float*, int)
     Write Thread (2, 0, 0) at 0x00000080 in /src/trap.cu:5193:Mag(float const *, float const *, int, float*, int)
     Current Value : 64, Incoming Value : 66

 WARN:(Warp Level Programming) Potential WAR hazard detected at __shared__ 0x307 in block (0, 0, 0) :
     Read Thread (0, 0, 0) at 0x000000a8 in /src/trap.cu:5197::Mag(float const *, float const *, int, float*, int)
     Write Thread (1, 0, 0) at 0x00000080 in /src/trap.cu:5193::Mag(float const *, float const *, int, float*, int)
     Current Value : 67, Incoming Value : 66

为什么不__synthreads(); 防止战争危害？

4

1 回答 1

2

当代码执行第 5197 和 5199 行然后循环回到第 5193 行时，没有干预 __synthreads()。

于 2013-09-27T20:57:56.167 回答