0

我想逐渐计算整个内核执行的线程执行数。是否有本机计数器或有任何其他方法可以做到这一点？我知道保持全局变量并由每个线程递增是行不通的，因为全局内存中的变量不能保证线程的同步访问。

4

3 回答 3

8

有许多方法可以衡量线程级执行效率。这个答案提供了不同收集机制的列表。Robert Crovella 的回答提供了一种手动检测方法，可以准确收集信息。可以使用类似的技术来收集内核中的散度信息。

为执行而启动的线程数（静态）

gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

启动的线程数

gridDim.x * gridDim.y * gridDim.z * ROUNDUP((blockDim.x * blockDim.y * blockDim.z), WARP_SIZE)

此数字包括在 warp 的生命周期内不活动的线程。

这可以使用 PM 计数器threads_launched 收集。

执行的 Warp 指令

计数器 inst_executed 计算执行/退出的 warp 指令的数量。

已发布经线指令

计数器 inst_issued 对发出的指令数进行计数。inst_issued >= inst_executed。某些指令将在每条指令执行时多次发出，以处理对狭窄执行单元的分派或为了处理共享内存和 L1 操作中的地址分歧。

执行的线程指令

计数器 thread_inst_executed 对执行的线程指令数进行计数。指标 avg_threads_executed_per_instruction 可以使用 thread_inst_executed / inst_executed 派生。此计数器的最大值是 WARP_SIZE。

不以执行的线程指令为前提

计算能力 2.0 及更高版本的设备使用指令预测来禁用扭曲中线程的回写，作为对短序列发散指令的性能优化。

计数器 not_predicated_off_thread_inst_executed 计算所有线程执行的指令数。此计数器仅适用于计算能力 3.0 及更高版本的设备。

not_predicated_off_thread_inst_executed <= thread_inst_executed <= WARP_SIZE * inst_executed

由于 thread_inst_executed 和 not_predicated_off_thread_inst_executed 计数器中的小错误，这种关系在某些芯片上会略微偏离。

分析器

Nsight Visual Studio Edition 2.x 支持收集上述计数器。

Nsight VSE 3.0 支持新的指令计数实验，可以收集每个 SASS 指令的统计信息，并以表格形式或在高级源、PTX 或 SASS 代码旁边显示数据。信息从 SASS 汇总到高级源。汇总的质量取决于编译器输出高质量符号信息的能力。建议您始终同时查看源代码和 SASS。该实验可以收集以下每条指令的统计信息：

一种。inst_executed b. thread_inst_executed（或活动掩码） c．not_predicated_off_thread_inst_executed（主动谓词掩码） d．active_mask e 的直方图。predicate_mask 的直方图

Visual Profiler 5.0 可以准确收集上述 SM 计数器。nvprof 可以收集和显示每个 SM 的详细信息。Visual Profiler 5.x 不支持收集 Nsight VSE 3.0 中可用的每条指令统计信息。旧版本的 Visual Profiler 和 CUDA 命令行分析器可以收集许多上述计数器，但结果可能不如 5.0 及更高版本的工具准确。

于 2013-05-29T02:08:25.433 回答
3

也许是这样的：

__global__ void mykernel(int *current_thread_count, ...){
  atomicAdd(current_thread_count, 1);
  // the rest of your kernel code
  }


int main() {
  int tally, *dev_tally;
  cudaMalloc((void **)&dev_tally, sizeof(int));
  tally = 0;
  cudaMemcpy(dev_tally, &tally, sizeof(int), cudaMemcpyHostToDevice);
  ....
  // set up block and grid dimensions, etc.
  dim3 grid(...);
  dim3 block(...)

  mykernel<<<grid, block>>>(dev_tally, ...);
  cudaMemcpy(&tally, dev_tally, sizeof(int), cudaMemcpyDeviceToHost); 
  printf("total number of threads that executed was: %d\n", tally);
  ....

  return 0;
  }

您可以在此处阅读有关原子函数的更多信息

许多人在评论中表达的困惑的部分原因是，当mykernel完成时（假设它运行成功），每个人都期望tally最终得到一个等于grid.x*grid.y*grid.z*block.x*block.y*block.z

于 2013-05-28T15:48:20.097 回答
0

我认为没有办法计算特定路径分支中的线程数。对于直方图，最好有以下内容：

PS：直方图是关于计算每种颜色的像素。

for (i=0; i<256; i++) // 256 colors, 1 pixel = 1 thread if (threadidx.x == i) Histogramme[i] = CUDA_NbActiveThreadsInBranch() // Threads having i as color

于 2015-03-17T18:40:16.753 回答