CUDA的调用图

Question

我正试图为我在埃及的CUDA中的代码生成调用图，但通常的方法似乎不起作用（因为nvcc没有任何可以做与-fdump-rtl扩展相同的事情的标志）。

更多详细信息：

我有一个非常大的代码（我不是它的作者），它跨越了多个.cu文件，如果我有调用图，我会更容易理解它在做什么。

我敢打赌，这个问题的答案对其他人也有用。

关于如何使用cuda（.cu）文件做到这一点，有什么想法吗？

Answer 1

您可以通过clang 3.8的CUDA支持来实现这一点。

首先，编译CUDA代码以发出llvm（在安装了CUDA 7.5的Windows上的示例）：

clang++ -c main.cu --cuda-gpu-arch=sm_35 -o main.ll -I"C:Program FilesNVIDIA GPU Computing ToolkitCUDAv7.5include"

然后，使用生成的ll构建具有opt:的调用图

opt.exe main.ll -analyze -dot-callgraph

注意，opt不是默认二进制分发的一部分，您可能需要自己构建它（我有一个3.7.1版本，它已经能够从3.8管理ll）

示例main.cu文件：

#include <cuda_runtime.h>
__device__ int f() { return 1; }
__device__ float g(float* a) { return a[f()] ; }
__device__ float h() { return 42.0f ; }
__global__ void kernel (int a, float* b)
{
        int c = a + f();
        g(b);
        b[c] = h();
}

生成的点文件：

digraph "Call graph" {
        label="Call graph";
        Node0x1e3d438 [shape=record,label="{external node}"];
        Node0x1e3d438 -> Node0x1e3cfb0;
        Node0x1e3d438 -> Node0x1e3ce48;
        Node0x1e3d438 -> Node0x1e3d0a0;
        Node0x1e3d438 -> Node0x1e3d258;
        Node0x1e3d438 -> Node0x1e3cfd8;
        Node0x1e3d438 -> Node0x1e3ce98;
        Node0x1e3d438 -> Node0x1e3d000;
        Node0x1e3d438 -> Node0x1e3cee8;
        Node0x1e3d438 -> Node0x1e3d078;
        Node0x1e3d000 [shape=record,label="{__cuda_module_ctor}"];
        Node0x1e3d000 -> Node0x1e3ce98;
        Node0x1e3d000 -> Node0x1e3d168;
        Node0x1e3d078 [shape=record,label="{__cuda_module_dtor}"];
        Node0x1e3d078 -> Node0x1e3cee8;
        Node0x1e3cfb0 [shape=record,label="{^A?f@@YAHXZ}"];
        Node0x1e3d0a0 [shape=record,label="{^A?h@@YAMXZ}"];
        Node0x1e3ce48 [shape=record,label="{^A?g@@YAMPEAM@Z}"];
        Node0x1e3ce48 -> Node0x1e3cfb0;
        Node0x1e3d258 [shape=record,label="{^A?kernel@@YAXHPEAM@Z}"];
        Node0x1e3d258 -> Node0x1e3cfb0;
        Node0x1e3d258 -> Node0x1e3ce48;
        Node0x1e3d258 -> Node0x1e3d0a0;
        Node0x1e3d168 [shape=record,label="{__cuda_register_kernels}"];
        Node0x1e3cee8 [shape=record,label="{__cudaUnregisterFatBinary}"];
        Node0x1e3cee8 -> Node0x1e3d528;
        Node0x1e3cfd8 [shape=record,label="{__cudaRegisterFunction}"];
        Node0x1e3cfd8 -> Node0x1e3d528;
        Node0x1e3ce98 [shape=record,label="{__cudaRegisterFatBinary}"];
        Node0x1e3ce98 -> Node0x1e3d528;
}