scatter2.0.9 no matching function for call to '__ldg'

[zhong110020]

csrc/hip/segment_csr_cuda.hip:188:20: error: no matching function for call to '__ldg' scalar_t val = __ldg(src_data + row_idx); ^~~~~ csrc/hip/segment_csr_cuda.hip:275:27: note: in instantiation of function template specialization 'gather_csr_kernel<c10::Half, 4>' requested here hipLaunchKernelGGL(( gather_csr_kernel<scalar_t, 4>), dim3(BLOCKS(1, 4 * N)), dim3(THREADS), 0, stream,

[rusty1s]

Can you share some more information about your system and PyTorch version? Can you post the full installation log?

[zhong110020]

Can you share some more information about your system and PyTorch version? Can you post the full installation log?

pytorch11.8，use pytorch_scatter0.2.9version，add environment variables export FORCE_CUDA=1，Compile Command python3 setup.py install --user bdist_wheel，The compilation process reported an error as follows： ture date. Please use a non-deprecated interface with equivalent functionality instead. For a listing of replacement headers and interfaces, consult the file backward_warning.h. To disable this warning use -Wno-deprecated. [-W#warnings]

warning \

^ csrc/hip/segment_csr_cuda.hip:108:3: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] hipSetDevice(src.get_device()); ^~~~ ~~~~ csrc/hip/segment_csr_cuda.hip:228:3: warning: ignoring return value of function declared with 'nodiscard' attribute [-Wunused-result] hipSetDevice(src.get_device()); ^~~~ ~~~~ csrc/hip/segment_csr_cuda.hip:188:20: error: no matching function for call to 'ldg' scalar_t val = __ldg(src_data + row_idx); ^~~~~ csrc/hip/segment_csr_cuda.hip:275:27: note: in instantiation of function template specialization 'gather_csr_kernel<c10::Half, 4>' requested here hipLaunchKernelGGL(( gather_csr_kernel<scalar_t, 4>), dim3(BLOCKS(1, 4 N)), dim3(THREADS), 0, stream, ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:30:31: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const char *' for 1st argument device inline static char ldg(const char ptr) { return ptr; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:32:32: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const char2 ' (aka 'const HIP_vector_type<char, 2> ') for 1st argument device inline static char2 __ldg(const char2 ptr) { return ptr; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:34:32: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const char4 ' (aka 'const HIP_vector_type<char, 4> ') for 1st argument device inline static char4 ldg(const char4 ptr) { return ptr; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:36:38: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const signed char ' for 1st argument device inline static signed char ldg(const signed char ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:38:40: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const unsigned char ' for 1st argument device inline static unsigned char __ldg(const unsigned char ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:41:32: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const short ' for 1st argument device inline static short ldg(const short ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:43:33: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const short2 ' (aka 'const HIP_vector_type<short, 2> ') for 1st argument device inline static short2 ldg(const short2 ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:45:33: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const short4 ' (aka 'const HIP_vector_type<short, 4> ') for 1st argument device inline static short4 ldg(const short4 ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:47:41: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const unsigned short *' for 1st argument device inline static unsigned short ldg(const unsigned short ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:50:30: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const int ' for 1st argument device inline static int __ldg(const int ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:52:31: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const int2 ' (aka 'const HIP_vector_type<int, 2> ') for 1st argument device inline static int2 __ldg(const int2 ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:54:31: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const int4 ' (aka 'const HIP_vector_type<int, 4> ') for 1st argument device inline static int4 __ldg(const int4 ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:56:39: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const unsigned int ' for 1st argument device inline static unsigned int ldg(const unsigned int ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:59:31: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const long *' for 1st argument device inline static long ldg(const long ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:61:40: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const unsigned long ' for 1st argument device inline static unsigned long __ldg(const unsigned long ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:64:36: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const long long ' for 1st argument device inline static long long ldg(const long long ptr) { return ptr[0]; } ^ /opt/dtk/hip/include/hip/amd_detail/hip_ldg.h:66:36: note: candidate function not viable: no known conversion from 'const c10::Half ' to 'const longlong2 ' (aka 'const HIP_vector_type<long long, 2> ') for 1st argument device inline static longlong2 ldg(const longlong2* ptr) { return ptr[0]; }

[zhong110020]

cuda11.8，pytorch1.13

[rusty1s]

Can you try if running

pip install torch-scatter==2.0.9 --no-index -f https://data.pyg.org/whl/torch-1.13.0+cu117.html

fixes your issues?