// SPDX-License-Identifier: BSD-3-Clause
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#define CHECK(x) do { cudaError_t e_=(x); if(e_!=cudaSuccess){printf("CUDA_ERROR %s L%d\n", cudaGetErrorString(e_), __LINE__); printf("BLK512 DONE FAIL\n"); return 1;} } while(0)
__global__ void w512(int *out){ int i = (blockIdx.x*512 + threadIdx.x)*4;
  out[i]=i; out[i+1]=i+1; out[i+2]=i+2; out[i+3]=i+3; }
__global__ void w256(int *out){ int i = (blockIdx.y*16 + threadIdx.y)*256 + blockIdx.x*16 + threadIdx.x;
  out[i]=i; }
int main(void){
  int *d2, *d1; const int n=16384; int *h=(int*)malloc(n*4); int bad1=0,bad2=0,i;
  setvbuf(stdout,NULL,_IONBF,0);
  CHECK(cudaMalloc(&d2, n*4)); CHECK(cudaMalloc(&d1, n*4));
  CHECK(cudaMemset(d2,0,n*4)); CHECK(cudaMemset(d1,0,n*4));
  w512<<<8,512>>>(d2); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize());
  CHECK(cudaMemcpy(h,d2,n*4,cudaMemcpyDeviceToHost));
  for(i=0;i<n;i++) if(h[i]!=i) bad1++;
  printf("BLK512 w512 grid=(8,1) block=(512,1) bad=%d first8=%d %d %d %d %d %d %d %d\n",bad1,h[0],h[1],h[2],h[3],h[4],h[5],h[6],h[7]);
  CHECK(cudaMemset(d2,0,n*4));
  w256<<<dim3(16,4),dim3(16,16)>>>(d1); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize());
  CHECK(cudaMemcpy(h,d1,n*4,cudaMemcpyDeviceToHost));
  for(i=0;i<n;i++) if(h[i]!=i) bad2++;
  printf("BLK512 w256 grid=(16,4) block=(16,16) bad=%d first8=%d %d %d %d %d %d %d %d\n",bad2,h[0],h[1],h[2],h[3],h[4],h[5],h[6],h[7]);
  { int z; int firstbad=-1,lastbad=-1; for(z=0;z<n;z++){ if(h[z]!=z){ if(firstbad<0) firstbad=z; lastbad=z; } }
    printf("BLK512 w256 span firstbad=%d lastbad=%d | p0[1000..1023]=%d %d %d %d %d %d %d %d %d %d | p1[0..7]=%d %d %d %d %d %d %d %d | p3[1016..1023]=%d %d %d %d %d %d %d %d | p4[0..7]=%d %d %d %d %d %d %d %d | p15[1016..1023]=%d %d %d %d %d %d %d %d\n",
      firstbad,lastbad,
      h[1000],h[1001],h[1010],h[1015],h[1016],h[1017],h[1020],h[1021],h[1022],h[1023],
      h[1024],h[1025],h[1026],h[1027],h[1028],h[1029],h[1030],h[1031],
      h[4096-8],h[4096-7],h[4096-6],h[4096-5],h[4096-4],h[4096-3],h[4096-2],h[4096-1],
      h[4096],h[4097],h[4098],h[4099],h[4100],h[4101],h[4102],h[4103],
      h[16384-8],h[16384-7],h[16384-6],h[16384-5],h[16384-4],h[16384-3],h[16384-2],h[16384-1]); }
  { int b,bb; int zc=0,fc=0; int prevoff=0;
    for(b=0;b<256;b++){ int cb=0; for(bb=0;bb<64;bb++) if(h[b*64+bb]!=b*64+bb) cb++;
      if(cb==0) zc++; if(cb==64) fc++;
      if(b<16) printf("BLK512 w256 blk%d bad=%d\n",b,cb); }
    printf("BLK512 w256 blocks zero-bad=%d full-bad=%d\n",zc,fc); }

  printf("BLK512 DONE\n"); return 0;
}