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可以从特殊寄存器存储到内存吗？

回答

相关问题