shared memory 原子操作

加速计算 CUDA
1

下面是个 kernel函数，nisight调试出错，显示：error = access violation on load (shared memory)。在调用这个kernel的函数中，用cudaStatus = cudaDeviceSynchronize(); 检查，显示cudaDeviceSynchronize returned error code 30。
琢磨好久了，无果。请问，是什么个错误呢？
global static void ComputeMatrix(BYTE* NewImage, int LocalImageWidth,int LocalImageHeight,float final_sum1,int GPUPMH )
{
int xIndex=__mul24(blockDim.x,blockIdx.x)+threadIdx.x;
int yIndex=__mul24(blockDim.y,blockIdx.y)+threadIdx.y;
// int xIndex=blockDim.xblockIdx.x+threadIdx.x;
// int yIndex=blockDim.yblockIdx.y+threadIdx.y;
int tid_in_x=threadIdx.x;
int tid_in_y=threadIdx.y;
int tid_in_block=__mul24(threadIdx.y,blockDim.x)+threadIdx.x;
// int tid_in_block=threadIdx.yblockDim.x+threadIdx.x;
unsigned int index=yIndexLocalImageWidth+xIndex;
shared int sin[K][K]; //图像
shared int sgrayCoMatrixRD[G][G];//灰度共生矩阵
sin[tid_in_x][tid_in_y]=0;
__syncthreads();//块内线程同步
oMatrixRD[tid_in_x][tid_in_y]=0;
__syncthreads();//块内线程同步
//传值
sin[tid_in_x][tid_in_y]=(unsigned int)NewImage[index]; //听到这不动 sin 不变 tid _in_x 不变 都是0
__syncthreads();//块内线程同步
atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x-D][tid_in_y+D]][sin[tid_in_x][tid_in_y]],1);
atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x][tid_in_y]][sin[tid_in_x-D][tid_in_y+D]],1);
GPUPMH[index]=sgrayCoMatrixRD[tid_in_x][tid_in_y];
__syncthreads();//块内线程同步

}

2

LZ您好：

nsight已经提示您shared memory读取越界了，这就是您的问题所在。

3

您在shared memory上越界了，請根據右下角的NSight提示圖（每次調試的時候都會出現）定位到你被提示的行，修復此行所存在的BUG.

4

定位在原子操作 atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x-D][tid_in_y+D]][sin[tid_in_x][tid_in_y]],1);
上，我检查过，sin数组里的数，都小于G，怎么会越界呢？
tid_in_y+D越界了？这个怎么检查啊？

5

定位在原子操作 atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x-D][tid_in_y+D]][sin[tid_in_x][tid_in_y]],1);
上，我检查过，sin数组里的数，都小于G，怎么会越界呢？
tid_in_y+D越界了？这个怎么检查啊？

6

LZ您好：

是这样的，“sin[tid_in_x-D][tid_in_y+D]”这里将引起读取shared memory中的数组sin的时候越界。您这里没有给出D的值，但是如果D是一个正值的话，考虑tid_in_x==0和tid_in_y==0的情况，您需要访问sin[-D][D]，这显然已经向前越界了。

同时nsight给出的信息是读取shared memory越界，而不是写入越界，所以不是sgrayCoMatrixRD数组越界。

以及您这下应该更加熟悉使用nsight这个强大的工具了吧:)

祝您编码顺利~

7

先谢过了
我也曾经考虑过tid_in_x ==0的情况，

if((tid_in_x+D>0)&&(tid_in_x-D>0)&&(tid_in_y+D<G)&&(tid_in_x-D<G))
atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x-D][tid_in_y+D]][sin[tid_in_x][tid_in_y]],1);
atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x][tid_in_y]][sin[tid_in_x-D][tid_in_y+D]],1);

用过if 条件句后，用nsight 还是那个错误，我就去掉if 了，if 这样写错了？
还有什么办法可以判断 ，使 sin 内的参数在【0，G) 范围内吗？

8

（1）
你需要保证的多了，光弄这个无济于事：
sgrayCoMatrixRD[sin[tid_in_x-D][tid_in_y+D]][sin[tid_in_x][tid_in_y]]
可以分解成：
t0 = sin[tid_in_x-D][tid_in_y+D];
t1 = sin[tid_in_x][tid_in_y]
sgrayCoMatrixRD[t0][t1]
你要看到里面一共有4个下标都需要保证不越界。

（2）你犯了一个低级错误，你的if只能保护下面的第一条语句，
第二条语句照样可以越界。

请您考虑这两点。
我和ICE们将竭诚为您服务，但需要您认真考虑。

感谢您的深夜来访。

9

LZ您好：

您是要判断tid_in_x的范围么？
“if((tid_in_x+D>0)&&(tid_in_x-D>0)&&(tid_in_y+D<G)&&(tid_in_x-D<G))”这里面混入了一个tid_in_y。

10

哦 我翻了一下，那时候没写错
if((tid_in_x-D>0)&&(tid_in_y+D<G)&&(tid_in_x-D<G)&&(tid_in_y+D<G))
atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x-D][tid_in_y+D]][sin[tid_in_x][tid_in_y]],1);
atomicAdd(&sgrayCoMatrixRD[sin[tid_in_x][tid_in_y]][sin[tid_in_x-D][tid_in_y+D]],1);
判断 sin 里边的参数x y 方向都是[0，G)。

11

但是如8#玫瑰斑竹所言，你if忘记大括号了~