[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.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283963 91177308-0d34-0410-b5e6-96231b3b80d8