CUDA-Programming icon indicating copy to clipboard operation
CUDA-Programming copied to clipboard

P108 代码有误

Open yuchangminghit opened this issue 1 year ago • 2 comments

樊老师你好,第10章线程束基本函数与协作组中,如下这一段:

如果想要在循环内去掉对线程号的约束,又要避免出现读-写竞争,可以将相关代码 段改写如下: real v = s_y[tid]; for (int offset = 16; offset > 0; offset >>= 1) { v += s_y[tid + offset]; __syncwarp(); s_y[tid] = v; __syncwarp(); }

其中v += s_y[tid + offset];这一句可能出现访存越界,本地测试的报错为an illegal memory access was encountered,因为当前线程块大小为128,tid最大可能的取值为127,而s_y数组的大小为128,此时tid + offset值有可能大于127造成非法访存。这里把v += s_y[tid + offset];改为v+=s_y[tid+offset<blockDim.x?tid+offset:0];可以避免这个报错,但这样做就没有办法在循环内去掉对线程号的约束了。

不知理解是否有误,还请指教。

yuchangminghit avatar Feb 10 '25 09:02 yuchangminghit

非常感谢指出这个问题。我的写法确实欠考虑。这里确实会出现越界访问。感谢你一连指出两个bug。我会在主页列出。

brucefan1983 avatar Feb 10 '25 10:02 brucefan1983

我已经在主页列出了这个错误:

Image

brucefan1983 avatar Feb 10 '25 10:02 brucefan1983