0

在研究使用方式时，我正在摆弄一些 SASS %laneid。在浪费了某人生命的一分钟失态之后（对不起 - 你知道你是谁），我现在有以下几点：

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 无法从特殊寄存器中存储，还是错过了优化机会？

4

1 回答 1

2

您不能直接从特殊寄存器中存储；S2R读取特殊寄存器的值需要特殊操作 ( )。

基本原理：为访问特殊寄存器提供所有指令寻址模式将违背RISC理念，并且（以我的拙见）考虑到这种操作在实践中发生的可能性有多大，指令位不会被很好地使用。

于 2017-06-02T23:48:29.107 回答