integer arithmetic bug in CUDA 8.0.44 and 8.0.51

Makefile

NVCCFLAGS=-lineinfo
SM_FLAGS?=-gencode arch=compute_60,code=sm_60

all : works fails

works : reproduce.cu
	nvcc $< $(NVCCFLAGS) -o $@

fails : reproduce.cu
	nvcc $< $(NVCCFLAGS) $(SM_FLAGS) -o $@

run : works fails
	@echo this works
	@./works
	@echo this fails
	@./fails

clean :
	rm -vf ./works ./fails

reproduce.cu

#include <cuda.h>
#include <cstdio>

//s is expected to be 3 (if it is not, change the offset)
__global__ void cookie(int* d, int s, int offset = 1)
{
    int tid = blockIdx.x*blockDim.x + threadIdx.x;
    int begin = -s/2;
    int end =    s/2;

    for(int i = begin;i < (end + offset);++i)
        d[tid + i + offset] = i;
}


__global__ void keks(int s)
{
       int r = -s/2;
       printf("r= %d\n", r);
}

int main()
{
       keks<<<2,2>>> (3);
       cudaDeviceSynchronize();
       if(cudaGetLastError() != cudaSuccess){
           printf("Ouch, the keks burned!\n");
           return 1;
       }

       int *d_data = 0;
       int cardinal = 3;
       cudaMalloc(&d_data,4*cardinal*sizeof(int));
       cudaMemset(d_data,0,4*cardinal*sizeof(int));

       cookie<<<2,2>>> (d_data,cardinal);
       if(cudaGetLastError() != cudaSuccess){
           printf("Ouch, the cookie burned!\n");
           return 1;
       }

       int *h_data = (int*)malloc(4*cardinal*sizeof(int));
       cudaMemcpy(h_data,d_data,4*cardinal*sizeof(int),cudaMemcpyDeviceToHost);
       if(h_data[1]!=0){
           printf("Integer arithmetic on the GPU has a problem!\n");
           return 1;
       } 

       cudaFree(d_data);
       free(h_data);
       return 0;
}