We allocated the memory for CUdeviceptr for v1 object and type cast to uchar** type.

CUdeviceptr v1;
checkCudaResult(cuMemAlloc((CUdeviceptr*)&v1, 16));
checkCudaErrors(cudaMemset((void*)v1, 1, 16));
uchar** v1char = (uchar**)v1;
foo_double<<<4,4>>>(v1char);
checkCudaResult(cuMemFree(v1));

The above code throws an error of illegal memory access when v1char is accessed inside the kernel like the following,

_global__ void foo_double(uchar** val) {
  unsigned int tid = threadIdx_x;
  unsigned int bid = blockIdx_x * blockDim_x + tid;
  uchar* val_ = val[bid];
  printf("%d:%dn", bid,*val_);
}

No compilation or build error, is there any correct way to typecast v1{device pointer} to uchar** pointer and access the address correctly.