0

我有一个简单的 CUDA 内核，它计算一个非常大的字符串的 1000 字节片段中的 A 数。数据库的布局是为了合并内存访问。从内核返回后，我的 main 函数将设备数组复制results到主机上的一个以供进一步分析。

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char ch;
  if(id < jobs_todo) {
    for(i = 0; i < 1000; i += 1){
     ch = database[jobs_todo*i + id];
     if(ch == 'A') A++;
   }
  results[id] = A;
}

内核运行良好。但是，如果我results[id]=A用一些琐碎的东西 替换results[id]=10或只是注释掉该行，它会运行得更快（10 倍）并且使用更少的寄存器，如--ptxas-options=-v. 如果我注释掉该行，内核将无济于事。CUDA 编译器是否通过查看传递的参数知道这一点？所以它选择什么都不做？

4

1 回答 1

3

您看到的是编译器优化的结果。编译将修剪“死”代码，即不直接导致内存写入的代码。所以你的内核

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  float A = 0; int i; char ch;
  if(id < jobs_todo) {
    for(i = 0; i < 1000; i += 1){
     ch = database[jobs_todo*i + id];
     if(ch == 'A') A++;
   }
   results[id]=10;
}

被有效地优化为

__global__ void kernel(unsigned int jobs_todo, char* database, float* results ) {

  unsigned int id = threadIdx.x + blockIdx.x * blockDim.x;
  results[id]=10;
}

显然，缩减代码的寄存器占用空间和执行时间远低于完整代码。您可以通过将代码编译为 PTX 并检查发出的代码来确认这一点。

于 2013-01-01T20:32:14.693 回答