PTX ISA 之 BFS 代码分析

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21124049
// Cuda compilation tools, release 8.0, V8.0.44
// Based on LLVM 3.4svn
//

.version 5.0
.target sm_20
.address_size 64

    // .globl   _Z6KernelP4NodePiPbS2_S1_S2_i

.visible .entry _Z6KernelP4NodePiPbS2_S1_S2_i(
    .param .u64 _Z6KernelP4NodePiPbS2_S1_S2_i_param_0,
    .param .u64 _Z6KernelP4NodePiPbS2_S1_S2_i_param_1,
    .param .u64 _Z6KernelP4NodePiPbS2_S1_S2_i_param_2,
    .param .u64 _Z6KernelP4NodePiPbS2_S1_S2_i_param_3,
    .param .u64 _Z6KernelP4NodePiPbS2_S1_S2_i_param_4,
    .param .u64 _Z6KernelP4NodePiPbS2_S1_S2_i_param_5,
    .param .u32 _Z6KernelP4NodePiPbS2_S1_S2_i_param_6
)
{
    .reg .pred  %p<6>;
    .reg .b16   %rs<6>;
    .reg .b32   %r<26>;
    .reg .b64   %rd<33>;
//g_graph_mask[tid] == 1 表示未被访问过
//
    ld.param.u64    %rd10, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_0];
    ld.param.u64    %rd11, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_1];
    ld.param.u64    %rd12, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_2];
    ld.param.u64    %rd15, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_3];
    ld.param.u64    %rd13, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_4];
    ld.param.u64    %rd14, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_5];
    ld.param.u32    %r12, [_Z6KernelP4NodePiPbS2_S1_S2_i_param_6];
    cvta.to.global.u64  %rd1, %rd15;//visited p3
    mov.u32     %r13, %ctaid.x;
    shl.b32     %r14, %r13, 8;//8 = 2^8 or * 256
    mov.u32     %r15, %tid.x;
    add.s32     %r1, %r14, %r15;// global tid
    setp.ge.s32 %p1, %r1, %r12;//if(tid < no_of_nodes)
    @%p1 bra    BB0_7;

    cvt.s64.s32 %rd2, %r1;//rd2 是 global tid
    cvta.to.global.u64  %rd16, %rd12;
    add.s64     %rd3, %rd16, %rd2;
    ld.global.u8    %rs1, [%rd3]; // 根据 tid 来读取 g_graph_mask 的数据，Cache？ 无coherence 问题，可以cache
    setp.eq.s16 %p2, %rs1, 0;// g_graph_mask[tid] is 1
    @%p2 bra    BB0_7;

    cvta.to.global.u64  %rd17, %rd10;//g_graph_nodes first parameter
    mov.u16     %rs2, 0;
    st.global.u8    [%rd3], %rs2; // 赋值 为 0,表示 图中节点 被访问过了 ，之后要做判断，可cache
    add.s64     %rd18, %rd1, %rd2;
    mov.u16     %rs3, 1;
    st.global.u8    [%rd18], %rs3;// 赋值 为 1，表示 图中节点 被访问过了，之后要做判断，可cache
    shl.b64     %rd19, %rd2, 3;// gtid 为何左移 3位？
    add.s64     %rd20, %rd17, %rd19;//rd20 = rd17 : g_graph_nodes first parameter + gtid*8byte?
    add.s64     %rd4, %rd20, 4;//rd4 =? no_of_edages : Node second element
    ld.global.u32   %r19, [%rd20+4];//no_of_edages
    setp.lt.s32 %p3, %r19, 1;// for loop 的判断，是否有边？还是边的个数大于 1？
    @%p3 bra    BB0_7;

    cvta.to.global.u64  %rd21, %rd13;//rd13 gcost fifth parameter
    shl.b64     %rd22, %rd2, 2;// gtid 为何左移 2 位？4bytes?
    add.s64     %rd5, %rd21, %rd22;//rd5 = &gcost + gtid*4byte?
    ld.global.u32   %r24, [%rd4+-4];//read rd4:edages -4 -> starting ?  cache ?
    cvta.to.global.u64  %rd23, %rd11;//rd11 :parameter1 ->g_graph_edges
    mul.wide.s32    %rd24, %r24, 4;// node.starting 4bytes?
    add.s64     %rd32, %rd23, %rd24;// &starting+edges
    cvta.to.global.u64  %rd31, %rd14;// parameter5 -> gover
    mov.u32     %r23, %r24;//starting ?

BB0_4:
    mov.u32     %r21, %r24;//starting
    mov.u32     %r25, %r21;//starting
    ld.global.s32   %rd8, [%rd32];//id = read[i] : [staring + edges]  不能 预测
    add.s64     %rd25, %rd1, %rd8;//visited [0] +id
    ld.global.u8    %rs4, [%rd25];//read visited array Not cache      不能 预测
    setp.ne.s16 %p4, %rs4, 0;//is or not visited ?
    @%p4 bra    BB0_6;

    ld.global.u32   %r16, [%rd5];//read gcost cache depend gtid
    add.s32     %r17, %r16, 1;// gcost【tid】 +1
    shl.b64     %rd27, %rd8, 2;// *4 byte? int id?
    add.s64     %rd28, %rd21, %rd27;//*gcost 首地址 + id ×4 bytes
    st.global.u32   [%rd28], %r17;//gcost[id] = gcost[tid]+1; 与id 相关， 不可预测。 not cache ；

    add.s64     %rd30, %rd16, %rd8;
    st.global.u8    [%rd30], %rs3; // g_graph_mask[id]=true; not cache
    st.global.u8    [%rd31], %rs3;//*g_over=true; not cache
    ld.global.u32   %r25, [%rd4+-4];// read rd4 :edages -4 -> starting ?  cache ?

BB0_6:
    mov.u32     %r24, %r25;// &starting
    add.s64     %rd32, %rd32, 4;//&starting + 4 岂不是 no_of_edges
    add.s32     %r18, %r24, %r19;// starting + no_of_edges 
    add.s32     %r23, %r23, 1; // i++
    setp.lt.s32 %p5, %r23, %r18; // starting < starting + no_of_edges
    @%p5 bra    BB0_4;

BB0_7:
    ret;
}

程序流程图 bfs.sass