CUDA-Programming
CUDA-Programming copied to clipboard
P108 代码有误
樊老师你好,第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];可以避免这个报错,但这样做就没有办法在循环内去掉对线程号的约束了。
不知理解是否有误,还请指教。
非常感谢指出这个问题。我的写法确实欠考虑。这里确实会出现越界访问。感谢你一连指出两个bug。我会在主页列出。
我已经在主页列出了这个错误: