[CUDA] Add support for CUDA surfaces (llvm#132883)

This adds support for all the surface read and write calls to clang. It
extends the pattern used for textures to surfaces too.

I tested this by generating all the various permutations of the calls
and argument types in a python script, compiling them with both clang
and nvcc, and comparing the generated ptx for equivilence. They all
agree, ignoring register allocation, and some places where Clang picks
different memory write instructions. An example kernel is:

```
__global__ void testKernel(cudaSurfaceObject_t surfObj, int x, float2* result) {
    *result = surf1Dread<float2>(surfObj, x, cudaBoundaryModeZero);
}
```

---------

Signed-off-by: Austin Schuh <[email protected]>

Author: AustinSchuh

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

@@ -384,6 +384,7 @@ __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
 // will continue to fail as it does now.
 #endif // CUDA_VERSION
 #endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
+#include "surface_indirect_functions.h"
 #include "texture_fetch_functions.h"
 #include "texture_indirect_functions.h"