[CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.

[CUDA] Make touching a kernel from a __host__ __device__ function a deferred error.

Previously, this was an immediate, don't pass go, don't collect $200
error.  But this precludes us from writing code like

  __host__ __device__ void launch_kernel() {
    kernel<<<...>>>();
  }

Such code isn't wrong, following our notions of right and wrong in CUDA,
unless it's codegen'ed.

llvm-svn: 283963