Hello all, So I’ve come across a rather curious problem. When I call a device function that draws some random numbers and stores them in every second entry of a a shared array, then the array is not always displayed correctly when I print it, although I call __syncthreads() at the end of the functi…

So I’ve looked a bit further and it doesn’t seem to be a problem of the device function since if I put the code directly in the kernel, the behaviour doesn’t change, except for the fact that you need one less __syncthreads() before the fibonacci(int) function (which is a bit of a mystery to me, too)…

I’m not able to figure out what your claim is. I ran the code with and without CASE_B defined and in either case, it seemed that every other element output was 0.000000 If you want help, my suggestion would be: include additional information like platform, CUDA version, compile command, and GPU…

Hi txbob, thank you for your reply! The above program is designed to fill every entry of testFloat with 0., then fill every entry with even-numbered iterator with a random number. print every entry of testfloat #define CASE_A enables a __syncthreads() in the testFunction body #define CASE_B CAS…

It would appear to me that a __syncthreads() both before and after the if (threadIdx.x<128) block is a necessity. Christian

I commented out #define CASE_C and #define CASE_D and compiled and ran the code using CUDA 9.1 on Tesla P100 on CentOS 7. It appears that every odd-index output is zero. Perhaps you should try CUDA 9.1

@Christian , thank you for your reply. For me, this doesn’t change anything, unfortunately. I get what you mean and it should be put there but it changes nothing on my computer. @txbob, thanks for your reply. I tested the code with a GTX680 with cuda 6.5 on Ubuntu 14.04.3 LTS and got the same result…

Have you run this program through cuda-memcheck to see if there’s anything broken with respect to the random generator during the initialization or access of generator states?

I have now (thanks, great idea), though no errors came up in either case.

So I have tried to run this with CUDA 9.1 on my quadro card, as well as with the cooperative groups library (block.sync()) and it gave me the same behaviour. Still no errors in memcheck and still the same requirement for three consecutive threadfences (__syncthreads or block.sync).

Does it make a difference if you drop the random generators and write constant floating point values (e.g. 1.0f) into every 2nd element? Christian

Undefined Behaviour of syncthreads and branching inside a device__ function

Accelerated Computing CUDA CUDA Programming and Performance

cbuchner1 March 29, 2018, 2:48pm 12

Another thought: Wouldn’t it be better (safer) to have a random generator state per thread instead of per thread block?

Topic		Replies	Views
Bug report: Threads out of sync, branched syncthreads problem CUDA Programming and Performance	2	1703	November 30, 2009
Syncthreads and Stalling Kernels CUDA Programming and Performance	16	4143	August 26, 2010
cuda syncthreads fail CUDA Programming and Performance	7	3877	February 22, 2013
CUDA warp branching problem Stange bug due to unexpected warp branching order CUDA Programming and Performance	17	12218	July 11, 2010
Are __device__ functions with __syncthreads() a bad idea? CUDA Programming and Performance	2	97	September 18, 2024
Problems with __syncthreads() CUDA Programming and Performance	2	949	May 4, 2013
Why this kernel "needs" a __syncthreads() statement to work? CUDA Programming and Performance	3	897	May 7, 2017
problem with __syncthreads(); CUDA Programming and Performance	1	1696	December 15, 2011
__syncthreads screwes calculation CUDA Programming and Performance	2	3422	November 22, 2007
syncthreads error? CUDA Programming and Performance	16	33368	June 2, 2008

Undefined Behaviour of __syncthreads and branching inside a __device__ function

Related topics

Undefined Behaviour of syncthreads and branching inside a device__ function