学步园

 ERROR: Potential RAW hazard detected at __shared__：

========= ERROR: Potential RAW hazard detected at __shared__ 0xcf150 in block (2, 0, 0) :
=========     Write Thread (596, 0, 0) at 0x00001078 in track_ellipse_kernel.cu:167:IMGVF_kernel(float**, float**, int*, int*, float, float, float, int, float)
=========     Read Thread (38, 0, 0) at 0x00000748 in track_ellipse_kernel.cu:141:IMGVF_kernel(float**, float**, int*, int*, float, float, float, int, float)
=========     Current Value : 194
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/libcuda.so (cuLaunchKernel + 0x27b) [0x50b9b]
=========     Host Frame:/usr/local/cuda-5.5/lib/libcudart.so.5.5 [0xb7dc]
=========     Host Frame:/usr/local/cuda-5.5/lib/libcudart.so.5.5 (cudaLaunch + 0x140) [0x28c10]
=========     Host Frame:./leukocyte [0xc342]
=========

Error:  Priority level of report

RAW: Type of hazard

__shared__0xcf150(0x表示十六进制）:location of hazard

(2,0,0):block index(x,y,z)

Write, Read: access pattern

(596,0,0),(38,0,0): thread index

0x00001078:instruction offset in the kernel

track_ellipse_kernel.cu:167: file name and line number( if available)

IMGVF_kernel:kernel name

例子：

__global__ int kernel(void) {
int x;
__shared__ int a;
if (threadIdx.x == WRITER)
a = threadIdx.x;     //Write
x = a;               //Read
// do some work
}

Fixed Broadcast Implementation

__global__ int kernel(void) {
int x;
__shared__ int a;
if (threadIdx.x == WRITER)
a = threadIdx.x;
a
__syncthreads();
x = a;
// do some work
}

Stack Back Traces
 Saved host back trace at call site
— Precise errors : Kernel launch site
— Global Leaks : cudaMalloc site
— CUDA API errors : CUDA API call site
 Device function call back trace at error
 Supported host OS : Linux, Mac, Windows
 Supported devices : Fermi+
— Only in non blocking launch mode
 Enabled by default

Sample Back Trace
Invalid __local__ write of size 4
at 0x000000e8 in localRecursive.cu:24:recursive(int*)
by thread (6,0,0) in block (0,0,0)
Address 0x00fffbfc is out of bounds
Device Frame:recursive(int*) (fibonacci(int, int) : 0xe0)
Device Frame:recursive(int*) (fibonacci(int, int) : 0xe0)
Device Frame:recursive(int*) (fibonacci(int, int) : 0xe0)
Device Frame:recursive(int*) (recursive(int*) : 0x28)
Saved host backtrace up to driver entry point at kernel launch time
Host Frame:libcuda.so (cuLaunchKernel + 0x3ae) [0xcb8ae]
Host Frame:libcudart.so.5.0 [0x11dd4]
Host Frame:libcudart.so.5.0 (cudaLaunch + 0x182) [0x3ad82]
Host Frame:localRecursive (_Z28__device_stub__Z9recursivePiPi + 0x33) [0xfa3]
Host Frame:localRecursive (main + 0x2cd) [0x12ad]
Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
Host Frame:localRecursive [0xdc9]