nVIDIA GPU可以從特殊寄存器存儲到內存嗎？

我正在擺弄一些SASS，同時看着使用%laneid的方式。這浪費了別人的生命一分鐘（約抱歉 - 你知道你是誰）失態後，我現在有以下幾點：nVIDIA GPU可以從特殊寄存器存儲到內存嗎？

CUDA代碼：

__forceinline__ __device__ unsigned lane_id() 
{ 
    unsigned ret; 
    asm volatile ("mov.u32 %0, %laneid;" : "=r"(ret)); 
    return ret; 
} 

__global__ void dummy(unsigned *C) 
{ 
    C[0] = lane_id(); 
}

SASS（對SM 6.1）：

/*0008*/     MOV R1, c[0x0][0x20];   /* 0x4c98078000870001 */ 
    /*0010*/   {   MOV R2, c[0x0][0x140];   /* 0x4c98078005070002 */ 
    /*0018*/     S2R R0, SR_LANEID;  }  /* 0xf0c8000000070000 */ 
                   /* 0x001ffc011e2007ff */ 
    /*0028*/     MOV R3, c[0x0][0x144];   /* 0x4c98078005170003 */ 
    /*0030*/     STG.E [R2], R0;     /* 0xeedc200000070200 */ 
    /*0038*/     EXIT;       /* 0xe30000000007000f */ 
                   /* 0x001f8000fc0007ff */ 
    /*0048*/     BRA 0x40;      /* 0xe2400fffff07000f */ 
    /*0050*/     NOP;        /* 0x50b0000000070f00 */ 
    /*0058*/     NOP;        /* 0x50b0000000070f00 */ 
                   /* 0x001f8000fc0007e0 */ 
    /*0068*/     NOP;        /* 0x50b0000000070f00 */ 
    /*0070*/     NOP;        /* 0x50b0000000070f00 */ 
    /*0078*/     NOP;        /* 0x50b0000000070f00 */

因此，STG指令 - 商店全局內存我猜 - 不立即採取SR_LANEID，而是寄存器它被放置由內置PTX成。這是因爲（Pascal）GPU無法從特殊寄存器存儲，還是錯過了優化機會？

來源

2017-06-02 einpoklum

它需要一個特殊的操作（'S2R'）來讀取特殊寄存器。給所有指令尋址模式訪問特殊的寄存器將反擊[RISC]（https://en.wikipedia.org/wiki/Reduced_instruction_set_computer）的哲學和恕我直言不會被指令位花費，考慮到這種操作可能有多大在實踐中發生。 – tera

@tera：把它變成答案？ – einpoklum

您不能直接從專門的寄存器存儲;它需要一個特殊的操作（S2R）來讀取特殊寄存器的值。理由：給出所有指令尋址模式訪問特殊的寄存器將反擊RISC的哲學，並且（在我看來）不會花費指令位，因爲這種操作在實際中可能發生的可能性很大。

來源

2017-06-02 23:48:29 tera

你知道這個事實嗎？我的意思是，你不是在寫你認爲有意義的東西？ – einpoklum

我只寫了我認爲有意義的東西。我不知道這是事實，因爲SASS除了[指令名稱列表]外沒有記錄（http://docs.nvidia.com/cuda/cuda-binary-utilities/index.html#instruction-set-參考文獻），並附上一些描述性文字。 – tera

所以我+ 1'ed這個，但我不會接受之前，我有一些確證... – einpoklum

nVIDIA GPU可以從特殊寄存器存儲到內存嗎？

回答

相關問題