LLVM NVPTX 后端结构参数零大小

Posted 2023-02-16

【中文标题】LLVM NVPTX 后端结构参数零大小

【英文标题】：LLVM NVPTX backend struct parameter zero size

【发布时间】：2014-01-18 04:35:23

【问题描述】：

在加载由 LLVM 的 NVPTX 后端生成的 PTX 程序集时，我遇到了一个模糊的异常。 （我正在从 ManagedCuda 加载 PTX - http://managedcuda.codeplex.com/）

ErrorNoBinaryForGPU: This indicates that there is no kernel image available that is suitable for the device. This can occur when a user specifies code generation options for a particular CUDA source file that do not include the corresponding device configuration.

这是模块的 LLVM IR（有点奇怪，因为它是由工具生成的）

; ModuleID = 'Module'
target triple = "nvptx64-nvidia-cuda"

%testStruct = type  i32 

define void @kernel(i32 addrspace(1)*) 
entry:
  %1 = alloca %testStruct
  store %testStruct zeroinitializer, %testStruct* %1
  %2 = load %testStruct* %1
  call void @structtest(%testStruct %2)
  ret void


define void @structtest(%testStruct) 
entry:
  ret void


!nvvm.annotations = !!0

!0 = metadata !void (i32 addrspace(1)*)* @kernel, metadata !"kernel", i32 1

这是生成的 PTX

//
// Generated by LLVM NVPTX Back-End
//

.version 3.1
.target sm_20
.address_size 64

        // .globl       kernel
.visible .func structtest
(
        .param .b0 structtest_param_0
)
;

.visible .entry kernel(
        .param .u64 kernel_param_0
)

        .local .align 8 .b8     __local_depot0[8];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .s32       %r<2>;
        .reg .s64       %rl<2>;

        mov.u64         %rl1, __local_depot0;
        cvta.local.u64  %SP, %rl1;
        mov.u32         %r1, 0;
        st.u32  [%SP+0], %r1;
        // Callseq Start 0
        
        .reg .b32 temp_param_reg;
        // <end>
        .param .align 4 .b8 param0[4];
        st.param.b32    [param0+0], %r1;
        call.uni
        structtest,
        (
        param0
        );

        //
        // Callseq End 0
        ret;


        // .globl       structtest
.visible .func structtest(
        .param .b0 structtest_param_0
)



        ret;

我不知道如何阅读 PTX，但我感觉问题与 structtest 函数定义中 .param .b0 structtest_param_0 的 .b0 位有关。

传递非结构值（如整数或指针）可以正常工作，.b0.这样做时，函数的一部分会读取类似.b32 或.b64 之类的正常内容。

将三元组更改为 nvptx-nvidia-cuda（32 位）没有任何作用，包括/排除 http://llvm.org/docs/NVPTXUsage.html 中建议的数据布局

这是 NVPTX 后端的错误，还是我做错了什么？

---

更新：

我正在查看这个 - http://llvm.org/docs/doxygen/html/NVPTXAsmPrinter_8cpp_source.html - 它看起来好像该类型正在下降到行 01568，显然不是原始类型，并且 Ty->getPrimitiveSizeInBits() 返回零。 （至少这是我的猜测）

我是否需要添加一个特殊情况来检查它是否是一个结构，获取地址，生成参数byval，然后取消引用该结构？这似乎是一个 hacky 解决方案，但我不确定如何解决它。

【参考方案1】：

您是否尝试过从编译中获取错误消息缓冲区？在 managedCuda 中，这将类似于：

CudaContext ctx = new CudaContext();
CudaJitOptionCollection options = new CudaJitOptionCollection();
CudaJOErrorLogBuffer err = new CudaJOErrorLogBuffer(1024);
options.Add(err);
try

    ctx.LoadModulePTX("test.ptx", options);

catch 

    options.UpdateValues();
    MessageBox.Show(err.Value);                

当我运行你的 ptx 时，它会说：

ptxas 应用程序 ptx 输入，第 12 行；致命：'.b0'附近的解析错误：语法错误

ptxas fatal : Ptx 程序集因错误而中止”

什么支持你对 b0 的猜测。

【讨论】：

感谢您的信息！我不知道它存在，它肯定会在未来有所帮助。但是，它并没有真正解决我遇到的问题。