3
2
运行时
CUDA移植HIP编译错误
2023-06-25 10:43:13
146

转码一个CUDA 库(faiss)到HIP版本后,在编译时出现以下报错(ROCm 为4.5.2),这个错误出现在 Kernel 函数位置:

__global__ void ivfInterleavedScan(
                ^
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
/home/dengww/Ports/faiss/faiss/gpu/impl/IVFInterleaved.h:41:17: error: illegal SGPR to VGPR copy
fatal error: too many errors emitted, stopping now [-ferror-limit=]
clang-13: error: clang frontend command failed with exit code 70 (use -v to see invocation)
clang version 13.0.0 (/buildspace/das-build/zifang-4.5.2/llvm-project/clang 4fbbb57356b98000fde3816fd8b20aab5837974b)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-4.5.2/llvm/bin
clang-13: note: diagnostic msg: Error generating preprocessed source(s).

这里给出该 Kernel 函数的签名:

template <
        typename Codec,
        typename Metric,
        int ThreadsPerBlock,
        int NumWarpQ,
        int NumThreadQ,
        bool Residual>
__global__ void ivfInterleavedScan(
        Tensor<float, 2, true> queries,
        Tensor<float, 3, true> residualBase,
        Tensor<Index::idx_t, 2, true> listIds,
        void** allListData,
        int* listLengths,
        Codec codec,
        Metric metric,
        int k,
        // [query][probe][k]
        Tensor<float, 3, true> distanceOut,
        Tensor<int, 3, true> indicesOut);

在 Kernel 参数列表中,Tensor  是一个模板类:

template <
        typename T,
        int Dim,
        bool InnerContig = false,
        typename IndexT = int,
        template <typename U> class PtrTraits = traits::DefaultPtrTraits>
class Tensor;

从这个错误提示来看,大致是 Kernel 函数参数的问题(不太确定),非常感谢您的帮助!!!

全部评论(2)
只看楼主
T
tianlh
1楼
您好,方便提供完整的转码文件吗?我们这边调试看下
2023/6/26
一介散修大伯维
回复 @tianlh
已经解决了,问题是 CUDA PTX 代码的问题。
2023/7/21
没有更多内容了
--
楼主
Level.1
0 帖子
楼主热帖
暂无数据