0

我需要通过在不同场景下测试它们的吞吐量来测试一些 GPU。

这包括一个简单的 64b 乘法：

__device__ void add(unsigned int *data, bool flag){
unsigned int index = threadIdx.x;
unsigned int result;

asm ("{\n\t"
     "add.cc.u32    %1, %1, %1;\n\t"
     "addc.u32  %0, 0, 0;\n\t"
     "}"
     : "=r"(result), "+r"(index): );

if(flag)
    data[threadIdx.x] = result;
}

64b 模数：

__device__ void mod(){
    asm ("{\n\t"
         ".reg .u64 t1;\n\t"
         "cvt.u64.u32   t1, %0;\n\t"
         "rem.u64       t1, t1, t1;\n\t" 
         "}"
         : : "r"(index));
}

和 64b mul+mod：

__device__ void mulmod
    asm ("{\n\t"
         ".reg .u64 t1;\n\t"
         ".reg .u64 t2;\n\t"
         "mul.wide.u32  t1, %0, %0;\n\t"
         "cvt.u64.u32   t2, %0;\n\t"
         "rem.u64       t1, t1, t2;\n\t"
         "}"
         : : "r"(index));
}

我认为任何内存访问对我的意图都是完全没用的，我想使用线程索引变量作为输入。

而且由于我要在没有寄存器上写，所以我不需要关心寄存器的使用，我可以启动尽可能多的线程（每个 GPU 允许）

我想知道：

  • 如果这是正确的方法

  • 除了最大化线程数之外，是否有任何特定的线程配置可以检索到最佳吞吐量？

4

1 回答 1

2

您的第一个“子问题”的答案是否定的，这不是正确的方法，因为您编写的这些函数都不会被编译器发出。

您可以在我上面链接的问题的答案中看到更多详细信息，但简短的版本是 C 编译器级别的死代码优化将消除任何不参与写入内存的值的代码。因此，您必须让这些函数返回一个值，并且您必须以编译器无法推断出对您的设备函数的调用是多余的并消除它的方式使用返回值。

除此之外，每个 SM 必须有足够的活动扭曲来分摊架构中的所有指令调度延迟，并确保测量设备功能的指令吞吐量，而不是指令调度程序和流水线的延迟。

于 2013-02-03T19:12:26.063 回答