curand_init(*) weird error related to the size of curandState_t array

I am experiencing weird error, while I try to generate array of random numbers using cuRand.

Simplified code example:

__global__
void init_kernel_1D(const long long int size, const long long int seed, curandState_t * const od_state) {
    const long long int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        curandState_t l_state;
        curand_init(seed, idx, 0, &l_state);

        // Exit:
        od_state[idx] = l_state;
    }
}

template <typename T> __global__
void rand_kernel_1D(const long long int size, curandState_t * const od_state, T * const od_value) {
    const long long int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        curandState_t l_state = od_state[idx];
        T l_value = log(curand_normal_double(&l_state));

        // Exit:
        od_value[idx] = l_value;
    }
}

int main () {
    const long long int size = 300000;
    dim3 block (256, 1, 1);
    dim3 grid  (((size + block.x - 1) / block.x), 1, 1));

    // Pointers to device memory:
    curandState_t d_states; 
    double        d_values;

    // Allocate device memory:
    // ...

    // Start magic:
    init_kernel_1D <<<grid, block>>> (size, time(NULL), d_states);
    ERRCHK(cudaGetLastError());
    ERRCHK(cudaDeviceSynchronize());

    rand_kernel_1D <double> <<<grid, block>>> (size, d_states, d_values);
    ERRCHK(cudaGetLastError());
    ERRCHK(cudaDeviceSynchronize());
}

---

Few words of explanation:

1. In my real code, I am using smart pointer to device memory.
   Rreason: RAII.
   Functionality: Safe memory allocation and release, also data transfer between device and host.
   
2. ERRCHK is macro that "warps" template function located in separate compilation unit, and instantiated for 7 different error types (all cuda).
   Whenever error is detected this function throws an exception.

---

Error description:
Whenever I call init_kernel_1D <<<***>>> (***); while const long long int size is set to be above 200 000 (therefor d_states points to the array of 200 000 curandState_t elements located in device global memory) ERRCHK(cudaDeviceSynchronize()); reports an cudaErrorLaunchFailure error.

I haven't slightest idea why size of curandState_t array, causes this error.

1. I thought that I might be runing out of the device memory, but i have 4gb of ram on my gpu, so that's not a reason.
2. I double checked my array indexing method, but it works perfectly for every other data type. (tested on device array with size of: 300 000 000 elements and data type: double)