-2

I set clang++-12 as the CUDA compiler in CMake as follows:

cmake .. -DCMAKE_BUILD_TYPE=Debug \
  -DCMAKE_CUDA_ARCHITECTURES="75" \
  -DCMAKE_CUDA_COMPILER=clang++-12 \
  -DCMAKE_CXX_COMPILER=clang++-12 -DCMAKE_C_COMPILER=clang-12

I try to compile a large codebase with it, but essentially it fails for a .cu file that includes CUB and uses some textures with the following error:

clang: warning: Unknown CUDA version. cuda.h: CUDA_VERSION=11040. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]
clang: warning: argument unused during compilation: '-G -Werror' [-Wunused-command-line-argument]
In file included from /scratch/src/libgpu/src/skelapp/skel_dummy.cu:2:
In file included from /scratch/src/libgpu/include/gg.cuh:5:
In file included from /scratch/src/libgpu/include/AppendOnlyList.cuh:6:
In file included from /usr/local/cuda-11.4/include/cub/cub.cuh:61:
In file included from /usr/local/cuda-11.4/include/cub/device/device_spmv.cuh:41:
In file included from /usr/local/cuda-11.4/include/cub/device/dispatch/dispatch_spmv_orig.cuh:42:
In file included from /usr/local/cuda-11.4/include/cub/device/dispatch/../../agent/agent_spmv_orig.cuh:47:
/usr/local/cuda-11.4/include/cub/device/dispatch/../../agent/../iterator/tex_ref_input_iterator.cuh:133:39: error: reference to __host__ variable 'ref' in __device__ function
                words[i] = tex1Dfetch(ref, (tex_offset * TEXTURE_MULTIPLE) + i);
                                      ^
In file included from <built-in>:1:
In file included from /usr/lib/llvm-12/lib/clang/12.0.1/include/__clang_cuda_runtime_wrapper.h:333:
/usr/local/cuda-11.4/include/texture_indirect_functions.h:111:4: error: use of undeclared identifier '__nv_tex_surf_handler'
   __nv_tex_surf_handler("__itex1Dfetch", ptr, obj, x);
   ^
/usr/local/cuda-11.4/include/texture_indirect_functions.h:120:3: note: in instantiation of function template specialization 'tex1Dfetch<int>' requested here
  tex1Dfetch(&ret, texObject, x);
  ^
/scratch/src/libgpu/include/worklist.cuh:393:12: note: in instantiation of function template specialization 'tex1Dfetch<int>' requested here
    item = tex1Dfetch<int>(tx, id);
           ^
2 errors generated when compiling for sm_75.
make[2]: *** [libgpu/CMakeFiles/some_gpu.dir/build.make:146: libgpu/CMakeFiles/some_gpu.dir/src/skelapp/skel_dummy.cu.o] Error 1

How to work around this? Shall I submit bugs to NVIDIA or LLVM?

The environment is Ubuntu 20.04 with LLVM/Clang 12 and CUDA Toolkit 11.4.1 . CUB is bundled with CUDA Toolkit.

Serge Rogatch
  • 13,865
  • 7
  • 86
  • 158
  • 1
    Use a supported CUDA version, Clang only supports CUDA 9 and maybe CUDA 10, depending on how new your clang build is – talonmies Sep 05 '21 at 14:15
  • `clang: warning: Unknown CUDA version. cuda.h: CUDA_VERSION=11040. Assuming the latest supported version 10.1 [-Wunknown-cuda-version]` There would be no point submitting bugs to NVIDIA concerning this. NVIDIA doesn't develop or maintain Clang. A possible workaround is to use `nvcc` instead. – Robert Crovella Sep 05 '21 at 14:53
  • @RobertCrovella, the point is that `tex_ref_input_iterator.cuh:133` indeed contains a reference from a device function to the host-only variable `ref`, so CUB is broken and `nvcc` is broken too, therefore the error leaked unnoticed. – Serge Rogatch Sep 05 '21 at 16:15
  • 1
    That is a standard texture reference, perfectly legal since the first CUDA beta was released – talonmies Sep 05 '21 at 17:11
  • 2
    Clang does not support texture/surface lookups, regardless of the CUDA version used. It can be implemented, but there's been no specific need to justify the effort so far. – ArtemB Sep 07 '21 at 18:20
  • @ArtemB I think that would make a useful answer and would upvote such an answer. I'm imagining that applies to texture/surface *references* (what is in view here) and not objects. CUDA has deprecated texture/surface references, [recently](https://stackoverflow.com/questions/67193791/how-to-replace-the-deprecated-tex2dtexturet-2-cudareadmodeelementtype-floa). – Robert Crovella Sep 11 '21 at 17:26
  • 1
    FYI, clang may finally grow some support for texture lookups, soon. https://reviews.llvm.org/D110089 – ArtemB Sep 20 '21 at 18:36
  • @RobertCrovella -- clang has to deal with multiple CUDA versions, so we need to support both references and objecs. AFACIT, the distinction is largely user-facing. Under the hood both translate into the same code that passes some sort of opaque texture handle in a 64-bit register. So, to answer your question, both access methods are not supported now and both should be working for the foreseeable future once they are implemented (modulo what can be accessed by the user via the headers of the specific CUDA SDK version). – ArtemB Sep 20 '21 at 18:56

1 Answers1

3

How to work around this? Shall I submit bugs to NVIDIA or LLVM?

There's no workaround. Right now Clang does not support texture/surface lookups, regardless of the CUDA version used. There's already a bug open: https://bugs.llvm.org/show_bug.cgi?id=26400

Texture support can be implemented, but there's been no specific need to justify the effort so far.

ArtemB
  • 3,496
  • 17
  • 18