New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Reduce last warp, the shared memory sdata shoule be set volatile #3

Open

hitblackjack opened this issue Dec 30, 2019 · 0 comments

hitblackjack commented Dec 30, 2019

DeviceMatrix.cu
line 381
if(tid < warpSize){
for (size_t shift = warpSIze; shift>0;shift>>=1)
sdata[tid] += sdata[tid +shift];

without ”__syncthreads()“ here, the shared memory sdata should be set volatile.

reference:
https://stackoverflow.com/questions/21205471/cuda-in-warp-reduction-and-volatile-keyword

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment