Cuda error an illegal instruction was encountered

If I run the following program I get «an illegal instruction was encountered». The program accesses an array of unsigned long longs.

Error: ../file.cu(43) : CUDA Runtime API error 73: an illegal instruction was encountered.

GPU: GTX 760

Cuda 6.5 (V6.5.16). The error happens only when running the code in the Release mode (-O3 optimization flag).

The code:

#include <stdio.h>
#include <cuda_runtime.h>

typedef unsigned long long ul;

static inline void __checkCudaErrors(cudaError err, const char *file, const int line )
{
    if(cudaSuccess != err){
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.n",file, line, (int)err, cudaGetErrorString( err ) );
        exit(-1);
    }
}
#define checkCuda(err)  __checkCudaErrors (err, __FILE__, __LINE__)


__global__ void testReadOnly(ul *A)
{
    int i = threadIdx.x/32;
    A=A+(32*i);
    atomicCAS(A+threadIdx.x, 0ull, 0ull);
}
int main(void)
{
    ul *d_A = NULL;
    int size = 32*sizeof(ul);
    checkCuda(cudaMalloc((void **)&d_A, size));
    checkCuda(cudaMemset(d_A, (int)0, size));
    testReadOnly<<<1, 32>>>(d_A);
    checkCuda(cudaDeviceSynchronize());
    checkCuda(cudaFree(d_A));
    checkCuda(cudaDeviceReset());
    return 0;
}

PTX code: _Z12testReadOnlyPy:

MOV R1, c[0x0][0x44]
S2R R0, SR_TID.X
MOV R4, c[0x0][0x140]
LOP.AND R3, R0, -0x20
MOV R8, RZ
IADD R2.CC, R3, R0
ISET.LT.AND R0, R3, RZ, PT

IADD.X R0, R0, RZ
SHL R3, R0, 0x3
MOV R9, RZ
IMAD.U32.U32 R4.CC, R2, 0x8, R4
IMAD.U32.U32.HI R0, R2, 0x8, R3
IADD.X R5, R0, c[0x0][0x144]
ATOM.E.CAS.U64 RZ, [R4], R8, RZ

cudaError_t 73 : «an illegal instruction was encountered» occurs sometimes.

When pre-training the Dnn model(using nnet), the first 2 RBM layers run OK, then goto RBM 3 layers, this error occurs. If restart the training, this error may disappear and may occcur when goto RBM 4 layers.

RBM Log:

rbm-train-cd1-frmshuff —learn-rate=0.4 —l2-penalty=0.0002 —num-iters=20 —drop-data=0.0 —verbose=1 ‘—feature-transform=nnet-concat exp/dnn4_pretrain-dbn/tr_splice5-1_cmvn-g.nnet exp/dnn4_pretrain-dbn/2.dbn — |’ exp/dnn4_pretrain-dbn/3.rbm.init ‘ark:copy-feats scp:exp/dnn4_pretrain-dbn/train.scp ark:- |’ exp/dnn4_pretrain-dbn/3.rbm
LOG (rbm-train-cd1-frmshuff:SelectGpuIdAuto():cu-device.cc:241) Selecting from 1 GPUs
LOG (rbm-train-cd1-frmshuff:SelectGpuIdAuto():cu-device.cc:256) cudaSetDevice(0): GeForce GTX TITAN Black free:6055M, used:88M, total:6143M, free/total:0.985606
…
LOG (rbm-train-cd1-frmshuff:Init():nnet-randomizer.cc:31) Seeding by srand with : 777
LOG (rbm-train-cd1-frmshuff:main():rbm-train-cd1-frmshuff.cc:141) RBM TRAINING STARTED
LOG (rbm-train-cd1-frmshuff:main():rbm-train-cd1-frmshuff.cc:144) Iteration 1/20
VLOG[1] (rbm-train-cd1-frmshuff:main():rbm-train-cd1-frmshuff.cc:246) Setting momentum 0.9 and learning rate 0.2 after processing 0.000277778h

ERROR (rbm-train-cd1-frmshuff:SetZero():cu-matrix.cc:444) cudaError_t 73 : «an illegal instruction was encountered» returned from ‘cudaMemset2D(data_, stride_ * sizeof(Real), 0, num_cols_ * sizeof(Real), num_rows_ )’

ERROR (rbm-train-cd1-frmshuff:Free():cu-device.cc:654) cudaError_t 73 : «an illegal instruction was encountered» returned from ‘cudaFree(addr)’

[CUDA_Runtime_API.pdf]

cudaErrorIllegalInstruction = 73
The device encountered an illegal instruction during kernel execution The context
cannot be used, so it must be destroyed (and a new one should be created). All
existing device memory allocations from this context are invalid and must be
reconstructed if the program is to continue using CUDA.

Thanks in advance！

Feiteng