[clang] 8173405 - [CUDA] make use of deprecated texture API conditional on CUDA version.

[Artem Belevich via cfe-commits]

Author: Artem Belevich
Date: 2022-11-17T11:38:16-08:00
New Revision: 817340569bf98b696329c53508a0d87cc0daec25

URL: https://github.com/llvm/llvm-project/commit/817340569bf98b696329c53508a0d87cc0daec25
DIFF: https://github.com/llvm/llvm-project/commit/817340569bf98b696329c53508a0d87cc0daec25.diff

LOG: [CUDA] make use of deprecated texture API conditional on CUDA version.

Added: 
    

Modified: 
    clang/lib/Headers/__clang_cuda_texture_intrinsics.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_cuda_texture_intrinsics.h b/clang/lib/Headers/__clang_cuda_texture_intrinsics.h
index 3c0f0026f1f02..a71952211237b 100644
--- a/clang/lib/Headers/__clang_cuda_texture_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_texture_intrinsics.h
@@ -666,6 +666,7 @@ __device__ static void __tex_fetch(__T *__ptr, cudaTextureObject_t __handle,
       __tex_fetch_v4<__op>::template __run<__FetchT>(__handle, __args...));
 }
 
+#if CUDA_VERSION < 12000
 // texture<> objects get magically converted into a texture reference.  However,
 // there's no way to convert them to cudaTextureObject_t on C++ level. So, we
 // cheat a bit and use inline assembly to do it. It costs us an extra register
@@ -713,6 +714,7 @@ __tex_fetch(__DataT *, __RetT *__ptr,
       __tex_fetch_v4<__op>::template __run<__FetchT>(
           __tex_handle_to_obj(__handle), __args...));
 }
+#endif // CUDA_VERSION
 } // namespace __cuda_tex
 } // namespace
 #pragma pop_macro("__ASM_OUT")