CUDA __global__ function 参数分析

在论坛里面讨论到一个问题,__global__函数里面传递的参数,到底是怎么传输到每一个thread的,然后做了以下的一些分析;

这个是问题讨论帖子:http://topic.csdn.net/u/20090210/22/2d9ac353-9606-4fa3-9dee-9d41d7fb2b40.html

 

C/C++ code
__global__ static void HelloCUDA(char* result, int num)
{
    __shared__ int i;
    i = 0;
    char p_HelloCUDA[] = "Hello CUDA!";
    for(i = 0; i < num; i++) {
        result[i] = p_HelloCUDA[i];
    }
}

 
 
PTX code 
    .const .align 1 .b8 __constant432[12] = {0x48,0x65,0x6c,0x6c,0x6f,0x20,0x43,0x55,0x44,0x41,0x21,0x0};

    .entry _Z9HelloCUDAPci
    {
    .reg .u16 %rh<3>;
    .reg .u32 %r<16>;
    .reg .pred %p<4>;
    .param .u32 __cudaparm__Z9HelloCUDAPci_result;
    .param .s32 __cudaparm__Z9HelloCUDAPci_num;
    .local .align 4 .b8 __cuda___cuda_p_HelloCUDA_168[12];
    .shared .s32 i;
    .loc    14    15    0
$LBB1__Z9HelloCUDAPci:
    mov.u32     %r1, __constant432;      //
    mov.u32     %r2, __cuda___cuda_p_HelloCUDA_168;    //
    ld.const.u32     %r3, [%r1+0];       // id:17 not_variable+0x0
    st.local.u32     [%r2+0], %r3;       // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
    ld.const.u32     %r4, [%r1+4];       // id:17 not_variable+0x0
    st.local.u32     [%r2+4], %r4;       // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
    ld.const.u32     %r5, [%r1+8];       // id:17 not_variable+0x0
    st.local.u32     [%r2+8], %r5;       // id:18 __cuda___cuda_p_HelloCUDA_168+0x0
    .loc    14    20    0
    mov.s32     %r6, 0;                  //
    ld.param.s32     %r7, [__cudaparm__Z9HelloCUDAPci_num];    // id:16 __cudaparm__Z9HelloCUDAPci_num+0x0
    mov.u32     %r8, 0;                  //
    setp.le.s32     %p1, %r7, %r8;       //
    @%p1 bra     $Lt_0_9;                //
    mov.s32     %r9, %r7;                //
    mov.u32     %r10, __cuda___cuda_p_HelloCUDA_168;    //
    mov.u32     %r11, __cuda___cuda_p_HelloCUDA_168;    //
    add.u32     %r12, %r7, %r11;         //
    ld.param.u32     %r13, [__cudaparm__Z9HelloCUDAPci_result];    // id:19 __cudaparm__Z9HelloCUDAPci_result+0x0
    mov.s32     %r14, %r9;               //
$Lt_0_7:
// Loop body line 20, nesting depth: 1, estimated iterations: unknown
    .loc    14    21    0
    ld.local.s8     %rh1, [%r10+0];      // id:20 __cuda___cuda_p_HelloCUDA_168+0x0
    st.global.s8     [%r13+0], %rh1;     // id:21
    add.u32     %r13, %r13, 1;           //
    add.u32     %r10, %r10, 1;           //
    setp.ne.s32     %p2, %r10, %r12;     //
    @%p2 bra     $Lt_0_7;                //
    st.shared.s32     [i], %r7;          // id:22 i+0x0
    bra.uni     $Lt_0_5;                 //
$Lt_0_9:
    st.shared.s32     [i], %r6;          // id:22 i+0x0
$Lt_0_5:
    .loc    14    23    0
    exit;                             //
$LDWend__Z9HelloCUDAPci:
    } // _Z9HelloCUDAPci
 
 
Cubin code
architecture {sm_10}
abiversion   {1}
modname      {cubin}
consts {
        name    = __constant432
        segname = const
        segnum  = 0
        offset  = 0
        bytes   = 12
    mem {
        0x6c6c6548 0x5543206f 0x00214144
    }
}
code {
    name = _Z9HelloCUDAPci
    lmem = 12
    smem = 28  // 我们注意这里的smem 的数量
    reg  = 3
    bar  = 0
    bincode {
        0x10000001 0x2400c780 0xd0000001 0x60c00780
        0x10000201 0x2400c780 0xd0000801 0x60c00780
        0x10000401 0x2400c780 0x307ccbfd 0x6c20c7c8
        0xd0001001 0x60c00780 0x10014003 0x00000280
        0x1000f801 0x0403c780 0x1000c805 0x0423c780
        0x00000005 0xc0000780 0xd4000009 0x40200780
        0x20018001 0x00000003 0xd00e0209 0xa0200780
        0x3000cbfd 0x6c2147c8 0x20018205 0x00000003
        0x1000a003 0x00000280 0x1000ca01 0x0423c780
        0x00000c01 0xe4200780 0x30000003 0x00000780
        0x00000c01 0xe43f0781
    }
}
 
 


这个是一段加了shared memory的也有constant 的ptx代码~还有cubin

 

从cubin来看,确实有可能是通过global memory进来的,不过一定是分发到各自的.param变量里面去的,因为在每一个thread 里面都是可以修改传进来的参数的;
从这点来看:大体应该是 参数( global memory--> (constant memory or shared memory) 然后broadcast to each thread--> 每一个thread都各自的register里面有了参数
大体应该是这样一个过程;

### TVM CUDA Compile 注册错误分析 当遇到 `Global PackedFunc tvm_callback_cuda_compile` 的注册检查失败问题时,通常是因为函数无法被正确加载或覆盖。此问题可能源于模块初始化阶段的冲突或其他运行环境配置不当。 #### 1. 错误原因解析 该错误的核心在于 `check_failed: can_override == false` 表明当前环境中已经存在同名的全局 `PackedFunc` 函数实例,并且不允许重复定义或覆盖[^1]。具体来说: - **TVM 模块机制**:TVM 使用动态库管理功能实现扩展支持(如 CUDA 编译)。如果多个版本的 TVM 或其依赖项共存于同一进程中,则可能导致命名空间污染。 - **CUDA 编译回调**:`tvm_callback_cuda_compile` 是用于处理 NVIDIA GPU 设备上代码编译逻辑的关键组件之一。它通过调用外部工具链完成 PTX 文件生成工作流。 #### 2. 解决方案建议 以下是几种常见的排查方法以及对应的修复措施: ##### 方法一:清理多余安装包 确认是否有不同版本或者未完全卸载干净的老版 TVM 库残留下来干扰正常操作流程。可以尝试重新构建项目源码并指定唯一路径来规避此类隐患: ```bash pip uninstall -y tvm git clone https://github.com/apache/tvm.git --recursive cd tvm && mkdir build && cd build cmake .. make -j$(nproc) export PYTHONPATH=$(pwd)/python:${PYTHONPATH} ``` ##### 方法二:调整插件优先级设置 有时为了兼容某些特定需求可能会手动引入额外的功能增强型 DLL/SO 文件,在这种情况下需要仔细审查这些附加资源是否会无意间篡改核心行为模式从而引发异常状况发生。修改相关参数控制加载顺序有助于缓解这一矛盾局面: 编辑 `.config/cmake/config.cmake` 添加如下选项前缀执行后续步骤: ```cmake set(USE_LLVM OFF CACHE BOOL "") set(USE_CUTLASS ON CACHE BOOL "") ``` ##### 方法三:强制允许覆写现有条目 虽然不推荐这样做因为容易埋下难以察觉的安全漏洞风险但是作为临时应急手段还是可行的——只需简单改动几行 Python 脚本即可达成目的: ```python from tvm import register_func, override_global_func override_global_func("tvm_callback_cuda_compile", new_function_definition, allow_override=True) ``` 以上三种方式各有优劣需根据实际场景灵活选用最合适的那一种实施补救行动直至彻底消除上述提示信息为止。 ### 示例代码片段展示如何自定义替换原有实现部分 下面给出一段简单的例子说明怎样创建一个新的替代品供系统采用而不破坏原始框架结构完整性: ```python import tvm from tvm.contrib import nvcc def my_custom_nvcc_compiler(code): """A customized NVCC compiler.""" flags = ["--generate-code=arch=compute_70,code=sm_70"] return nvcc.compile_cuda(code, target="ptx", options=flags) register_func("tvm_callback_cuda_compile", my_custom_nvcc_compiler, True) ```
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值