What is low bit ? curand_kernel.h Example from manual


I am now trying to fill my matrix with RNG (Random Number Generator).

Unfortunately, Host API can not be compiled, but device is ok.

Many question were poped up to my head.

  1. Page 14 state that “Difference seeds are guaranteed to produce different starting states and different sequences.”

    And then page 19 in the source code comment is "Each thread gets same seed, a different sequence number,

no offset". So offset and state are act like a seed. Am I right ?

  1. What does low bit mean to this code. I read wiki and still did not understand.
/* Check if low bit set */

if(x & 1) {

  1. How to fill this matrix d_2Darray; //Thanks to Sarnath
double **d_2Darray;

double *h_temp[N];

// Create 2D Array on device

// Do not destroy h_temp[]

        cudaMalloc((void **)&d_2Darray, N*sizeof(double *));

        for(i=0; i<N; i++)


           cudaMalloc( (void **)&h_temp[i], N*sizeof(double));


        cudaMemcpy(d_2Darray, h_temp, N*sizeof(double *), cudaMemcpyHostToDevice);

//Copy host** arrayn to device

        for(i=0; i<N; i++)


        cudaMemcpy(h_temp[i], A[i], N*sizeof(double), cudaMemcpyHostToDevice);


I use this style to access each element in matrix. One thread for one element.

__global__ void adda( double **A)


    int x = threadIdx.x + blockIdx.x * blockDim.x;

    int y = threadIdx.y + blockIdx.y * blockDim.y;

    //int offset = x + y * blockDim.x * gridDim.x;

    if (x < N && y < N){			

		A[x][y] = 44  ;



Can I replace 44 to be something like curand(&localState);

Any help, URL, although book would be appreciated.

if(x & 1)

tests whether the least significant bit in the binary representation of “x” is set. The numeric weight of the least significant bit is 1. Wikipedia has an article on the least significant bit here: http://en.wikipedia.org/wiki/Least_significant_bit

Thank you for your reply.

I am now trying to fill my matrix with random number. Still can not get it works.

curandState **devStates;

	float *h_temp2[N];

	/* Allocate space for prng states on device */

	CUDA_CALL(cudaMalloc((void **)&devStates, N *sizeof(curandState *)));


	   CUDA_CALL(cudaMalloc( (void **)&h_temp2[N], N*sizeof(curandState)));	


	CUDA_CALL(cudaMemcpy(devStates,	h_temp2,


				cudaMemcpyHostToDevice)); //HERE IS 109TH LINE.	


I got.

2Dcurand.cu(110): error: expected an expression

Any idea ?

My system is Ubuntu 10.10 with latest nvidia driver.

sarit@AH64D:~/AH64D$ nvcc --version

nvcc: NVIDIA (R) Cuda compiler driver

Copyright (c) 2005-2011 NVIDIA Corporation

Built on Thu_May_12_11:09:45_PDT_2011

Cuda compilation tools, release 4.0, V0.2.1221

If there any alternative way to fill my matrix with random number.

Do not hesitate to tell me.