Skip to content

Commit ccfb055

Browse files
committed
[CUDA] Implement experimental support for texture lookups.
The patch implements header-only support for testure lookups. The patch has been tested on a source file with all possible combinations of argument types supported by CUDA headers, compiled and verified that the generated instructions and their parameters match the code generated by NVCC. Unfortunately, compiling texture code requires CUDA headers and can't be tested in clang itself. The test will need to be added to the test-suite later. While generated code compiles and seems to match NVCC, I do not have any code that uses textures that I could test correctness of the implementation. Hence the experimental status. Differential Revision: https://reviews.llvm.org/D110089
1 parent 72dddce commit ccfb055

File tree

6 files changed

+810
-4
lines changed

6 files changed

+810
-4
lines changed

clang/lib/Headers/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,7 @@ set(files
4747
__clang_cuda_complex_builtins.h
4848
__clang_cuda_device_functions.h
4949
__clang_cuda_intrinsics.h
50+
__clang_cuda_texture_intrinsics.h
5051
__clang_cuda_libdevice_declares.h
5152
__clang_cuda_math_forward_declares.h
5253
__clang_cuda_runtime_wrapper.h

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

Lines changed: 30 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -64,9 +64,9 @@
6464
#endif
6565

6666
// Make largest subset of device functions available during host
67-
// compilation -- SM_35 for the time being.
67+
// compilation.
6868
#ifndef __CUDA_ARCH__
69-
#define __CUDA_ARCH__ 350
69+
#define __CUDA_ARCH__ 9999
7070
#endif
7171

7272
#include "__clang_cuda_builtin_vars.h"
@@ -330,6 +330,34 @@ static inline __device__ void __brkpt(int __c) { __brkpt(); }
330330

331331
#pragma pop_macro("__host__")
332332

333+
// __clang_cuda_texture_intrinsics.h must be included first in order to provide
334+
// implementation for __nv_tex_surf_handler that CUDA's headers depend on.
335+
// The implementation requires c++11 and only works with CUDA-9 or newer.
336+
#if __cplusplus >= 201103L && CUDA_VERSION >= 9000
337+
// clang-format off
338+
#include <__clang_cuda_texture_intrinsics.h>
339+
// clang-format on
340+
#else
341+
#if CUDA_VERSION >= 9000
342+
// Provide a hint that texture support needs C++11.
343+
template <typename T> struct __nv_tex_needs_cxx11 {
344+
const static bool value = false;
345+
};
346+
template <class T>
347+
__host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
348+
cudaTextureObject_t obj,
349+
float x) {
350+
_Static_assert(__nv_tex_needs_cxx11<T>::value,
351+
"Texture support requires C++11");
352+
}
353+
#else
354+
// Textures in CUDA-8 and older are not supported by clang.There's no
355+
// convenient way to intercept texture use in these versions, so we can't
356+
// produce a meaningful error. The source code that attempts to use textures
357+
// will continue to fail as it does now.
358+
#endif // CUDA_VERSION
359+
#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
360+
#include "texture_fetch_functions.h"
333361
#include "texture_indirect_functions.h"
334362

335363
// Restore state of __CUDA_ARCH__ and __THROW we had on entry.

0 commit comments

Comments
 (0)