Windows平台下HIP/HIPRT应用程序调试方法咨询

1. 核函数直接打印输出 #

用HIP支持的printf()在核函数中输出变量，配合hipDeviceSynchronize()确保设备端输出同步到宿主控制台，快速验证变量状态：

__global__ void traceKernel() {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    float sampleValue = calculateSample(tid);
    printf("Thread %d: sample = %.4f\n", tid, sampleValue);
}

int main() {
    traceKernel<<<8, 64>>>();
    hipDeviceSynchronize(); // 必须同步才能捕获设备端printf输出
    return 0;
}

2. 使用hip-gdb调试（WSL2环境） #

通过WSL2安装ROCm环境，将HIP代码迁移或交叉编译到WSL2中运行，用hip-gdb实现核函数的逐行调试：

• 启动调试：hip-gdb ./path-tracer-exe
• 设置核函数断点：break traceKernel
• 启动程序：run
• 逐行步进：step，查看变量：print sampleValue
  该工具支持断点设置、寄存器/内存查看，功能对标CUDA的Nsight。

3. 宿主端校验核函数输出 #

将核函数中的关键数据写入全局内存，执行完成后拷贝到宿主端打印或写入日志，适合批量线程的状态排查：

__global__ void traceKernel(float* d_output) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    d_output[tid] = finalRayColor(tid);
}

int main() {
    const int threadCount = 512;
    float* d_output;
    hipMalloc(&d_output, threadCount * sizeof(float));
    
    traceKernel<<<8, 64>>>(d_output);
    hipDeviceSynchronize();
    
    float* h_output = (float*)malloc(threadCount * sizeof(float));
    hipMemcpy(h_output, d_output, threadCount * sizeof(float), hipMemcpyDeviceToHost);
    
    // 打印或写入日志
    for (int i = 0; i < threadCount; i++) {
        fprintf(stderr, "Thread %d: color = %.4f\n", i, h_output[i]);
    }
    
    free(h_output);
    hipFree(d_output);
    return 0;
}

4. 强化HIP API错误检查 #

封装错误检查宏，快速定位宿主端的内存分配、核函数启动等错误：

#define HIP_CHECK(err) \
    do { \
        hipError_t errCode = err; \
        if (errCode != hipSuccess) { \
            fprintf(stderr, "HIP Error: %s (%d) at %s:%d\n", \
                hipGetErrorString(errCode), errCode, __FILE__, __LINE__); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

// 使用示例
HIP_CHECK(hipMalloc(&d_sceneData, sizeof(SceneData)));
HIP_CHECK(hipLaunchKernelGGL(traceKernel, dim3(8), dim3(64), 0, 0, d_sceneData, d_output));

5. VS Code搭配ROCm插件调试 #

在VS Code中安装ROCm Debugger插件，通过WSL2连接到Linux环境，实现类似Visual Studio的图形化调试体验：支持核函数断点、变量实时查看、步进执行，操作逻辑和Nsight接近。

内容的提问来源于stack exchange，提问作者Tom Clabault