几乎相同的 CUDA 内核的低处理器效率
Low processor efficiency with almost same CUDA kernels
我创建了三个合成 CUDA 内核，它们几乎都只进行算术运算。所有三个内核都是相同的，只是它们中的每一个都执行不同数量的操作。内核 #1 执行 8 次操作，内核 #2 执行 16 次操作，内核 #3 执行 32 次。以下是这三者的 CUDA 内核实现。
内核#1：
#ifndef kernelWGSXMAPIXLLXOPS8_H_
#define kernelWGSXMAPIXLLXOPS8_H_
__global__ void WGSXMAPIXLLXOPS8 (const float *GIn, float *GOut, const float M, const float N, const float P) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float MF = (float) M;
float NF = (float) N;
float PF = (float) P;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
... // 8 FMA operations
temp1 = temp1 * MF + temp1;
GOut[gid] = temp1;
}
}
void WGSXMAPIXLLXOPS8_wrapper (const float *GIn, float *GOut,
const float M, const float N, const float P,
int numBlocks, int threadPerBlock) {
WGSXMAPIXLLXOPS8<<<numBlocks, threadPerBlock>>> (GIn, GOut, M, N, P);
}
#endif
内核#2：
#ifndef kernelWGSXMAPIXLLXOPS16_H_
#define kernelWGSXMAPIXLLXOPS16_H_
__global__ void WGSXMAPIXLLXOPS16 (const float *GIn, float *GOut, const float M, const float N, const float P) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float MF = (float) M;
float NF = (float) N;
float PF = (float) P;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
... // 16 FMA operations
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
GOut[gid] = temp1;
}
}
void WGSXMAPIXLLXOPS16_wrapper (const float *GIn, float *GOut,
const float M, const float N, const float P,
int numBlocks, int threadPerBlock) {
WGSXMAPIXLLXOPS16<<<numBlocks, threadPerBlock>>> (GIn, GOut, M, N, P);
}
#endif
内核 #3：
#ifndef kernelWGSXMAPIXLLXOPS32_H_
#define kernelWGSXMAPIXLLXOPS32_H_
__global__ void WGSXMAPIXLLXOPS32 (const float *GIn, float *GOut, const float M, const float N, const float P) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float MF = (float) M;
float NF = (float) N;
float PF = (float) P;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
... // 32 FMA operations
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
GOut[gid] = temp1;
}
}
void WGSXMAPIXLLXOPS32_wrapper (const float *GIn, float *GOut,
const float M, const float N, const float P,
int numBlocks, int threadPerBlock) {
WGSXMAPIXLLXOPS32<<<numBlocks, threadPerBlock>>> (GIn, GOut, M, N, P);
}
#endif
线程总数已设置为 16384，块大小为 256。我计算了每个内核的总 GFlops，分别为 20.44、56.53 和 110.12 GFlops。我试图想出一个解释，但什么也想不起来。所以我尝试使用 nvprof 并监控所有指标。所有指标几乎相等，以下是一些对我来说很重要的指标（我还包括内核 1 到 3 的结果）：
sm_efficiency_instance: 14.99, 16.78, 19.82 %
ipc_instance: 0.57 , 0.93 , 1.53
inst_replay_overhead: 0.399, 0.268, 0.165
dram_write_throughput: 18.08, 17.72, 16.9 GB/s
issued_ipc: 0.99 , 1.18 , 1.52
issue_slot_utilization: 19.48, 24.64, 33.76 %
stall_exec_dependency: 21.84, 26.38, 42.95 %
很明显，它们都具有相同的 dram_write_throughput，因为它们都向 DRAM 写入相同数量的数据，并且线程总数也相同。我不明白的是sm_efficiency。我的内核都在做算术（一样），怎么他们的sm_efficiency不一样。另外，为什么在同一个内核中有更多的算法会提高效率？我的理解是，他们都应该有同样的问题来寻找 warp 以定位在 SM 上。
任何人都可以使用以下指标帮助我了解 GFlops 的区别吗？
基本问题是您没有 "saturated" 可用的 GPU。有各种与内核启动相关的开销。如果内核花费在计算上的时间与这个开销相比很小，那么你的计算就会被开销所扭曲。
T = 开销时间(OT) + 计算时间(CT)
Flops/s = Flops/T = 翻牌/(OT + CT)
如果计算时间与开销时间相比较小（您的内核就是这种情况），那么您的计算将受到开销时间的影响。另一方面，如果与开销相比计算时间足够大，那么开销对结果的影响相对较小。
这是一个完整的测试案例，有几个案例运行，CUDA 9.1，Tesla P100 PCIE：
$ cat t79.cu
#ifndef SLEN
#define SLEN (8)
#endif
#ifndef NTPB
#define NTPB (256)
#endif
#ifndef BLKS
#define BLKS (16384/NTPB)
#endif
const size_t blks = BLKS;
const size_t ntpb = NTPB;
typedef float Ftype;
#include <iostream>
template <int LEN>
__global__ void WGSXMAPIXLLXOPS (Ftype *GOut, const Ftype M) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
Ftype MF = (Ftype) M;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
if (LEN > 8){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 16){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 32){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 64){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 128){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 256){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
#ifdef NO_WRITE
if (temp1 == -10.0)
#endif
GOut[gid] = temp1;
}
}
int main(){
float et;
Ftype *GOut;
const Ftype M = 1.0;
cudaMalloc(&GOut, blks*ntpb*sizeof(Ftype));
cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
WGSXMAPIXLLXOPS<SLEN><<<blks, ntpb>>> (GOut, M);
cudaDeviceSynchronize();
cudaEventRecord(start);
WGSXMAPIXLLXOPS<SLEN><<<blks, ntpb>>> (GOut, M);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&et, start, stop);
unsigned long long flpcnt = SLEN*2*blks*ntpb;
float Kflops_s = flpcnt/et;
std::cout << "MFlops per sec: " << Kflops_s/1000 << " kernel time: " << et << "ms" << std::endl;
cudaDeviceSynchronize();
}
$ nvcc -arch=sm_60 -o t79 t79.cu
$ ./t79
MFlops per sec: 14371.9 kernel time: 0.01824ms
$ nvprof ./t79
==14676== NVPROF is profiling process 14676, command: ./t79
MFlops per sec: 10101.1 kernel time: 0.025952ms
==14676== Profiling application: ./t79
==14676== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 100.00% 3.2320us 2 1.6160us 1.2480us 1.9840us void WGSXMAPIXLLXOPS<int=8>(float*, float)
API calls: 98.31% 389.62ms 1 389.62ms 389.62ms 389.62ms cudaMalloc
1.10% 4.3574ms 376 11.588us 357ns 465.31us cuDeviceGetAttribute
0.42% 1.6829ms 4 420.73us 272.19us 642.45us cuDeviceTotalMem
0.12% 487.27us 4 121.82us 90.094us 164.09us cuDeviceGetName
0.02% 80.363us 2 40.181us 15.789us 64.574us cudaLaunch
0.00% 17.118us 2 8.5590us 8.1400us 8.9780us cudaDeviceSynchronize
0.00% 13.118us 2 6.5590us 5.4290us 7.6890us cudaEventRecord
0.00% 10.603us 2 5.3010us 1.2440us 9.3590us cudaEventCreate
0.00% 8.5080us 8 1.0630us 460ns 1.7500us cuDeviceGet
0.00% 8.4590us 1 8.4590us 8.4590us 8.4590us cudaEventElapsedTime
0.00% 7.1350us 1 7.1350us 7.1350us 7.1350us cudaEventSynchronize
0.00% 6.8430us 4 1.7100us 180ns 5.9720us cudaSetupArgument
0.00% 4.7800us 3 1.5930us 437ns 2.8480us cuDeviceGetCount
0.00% 2.3490us 2 1.1740us 361ns 1.9880us cudaConfigureCall
$ nvcc -arch=sm_60 -o t79 t79.cu -DSLEN=512 -DBLKS=32768 -DNTPB=1024
$ ./t79
MFlops per sec: 8.08072e+06 kernel time: 4.25206ms
$
$ nvprof --metrics sm_efficiency_instance,ipc_instance,issued_ipc,issue_slot_utilization,stall_exec_dependency ./t79
==15447== NVPROF is profiling process 15447, command: ./t79
==15447== Some kernel(s) will be replayed on device 0 in order to collect all events/metrics.
Replaying kernel "void WGSXMAPIXLLXOPS<int=512>(float*, float)" (done)
Replaying kernel "void WGSXMAPIXLLXOPS<int=512>(float*, float)" (done)
MFlops per sec: 193432 kernel time: 177.632ms
==15447== Profiling application: ./t79
==15447== Profiling result:
==15447== Metric result:
Invocations Metric Name Metric Description Min Max Avg
Device "Tesla P100-PCIE-16GB (0)"
Kernel: void WGSXMAPIXLLXOPS<int=512>(float*, float)
2 issued_ipc Issued IPC 1.972106 1.972388 1.972247
2 issue_slot_utilization Issue Slot Utilization 98.23% 98.24% 98.24%
2 stall_exec_dependency Issue Stall Reasons (Execution Dependency) 16.35% 16.36% 16.36%
2 ipc Executed IPC 1.971976 1.972254 1.972115
2 sm_efficiency Multiprocessor Activity 99.78% 99.78% 99.78%
$
第一个 运行，数字与您的匹配（16384 个线程，每个块 256 个线程，8 条 FFMA 指令）显示内核持续时间约为 17us。然而，当我们在分析器中 运行 这种情况时，我们观察到实际内核执行只有大约 1.5us，剩下的是各种开销，包括内核启动延迟，以及使用 cudaEvent系统进行计时。所以这会使数字偏离。
另一方面，当我们启动大量块、每个块的线程数以及每个线程工作时，我们得到的数字是 P100 峰值能力的 80%。
当您从内核 1 到内核 3 时，您的大多数指标都在增加（变得更好）（dram 吞吐量除外，这是合理的。随着内核时间的增加，对于相同数量的写入数据，dram 平均吞吐量下跌降落）。这与为 GPU 提供更多工作是一致的，因此它可以隐藏各种延迟并分摊大量工作的开销。
让我们看一下上述最终 运行/"large" 内核的一些指标：
2 issued_ipc Issued IPC 1.972106 1.972388 1.972247
2 issue_slot_utilization Issue Slot Utilization 98.23% 98.24% 98.24%
2 stall_exec_dependency Issue Stall Reasons (Execution Dependency) 16.35% 16.36% 16.36%
2 ipc Executed IPC 1.971976 1.972254 1.972115
2 sm_efficiency Multiprocessor Activity 99.78% 99.78% 99.78%
IPC 大约为每个时钟 2，高于您的内核 3。请注意，此处的 IPC 为 2 是一个合理的上限：sm_60 SM 有 64 single-precision 个单元，足以安排每个时钟 2 条 FFMA 指令。
SM 效率和 issue_slot_utilization 是相似的指标。这意味着大约 98% 的时间，SM 可以在任何给定的时钟周期内发出一条或多条指令。
stall（exec 依赖项）正在回答问题，"across all the actual stall situations, what percent were due to execution dependency?"。您的内核在每行源代码之间具有执行依赖性——因为每行都依赖于前一行的结果。这意味着在汇编级别，每条 FFMA 指令都将取决于前一条指令的结果，因此只有在前一条指令完成后才能发出。
如果 SM 的可用工作订阅不足，那么 stall exec 依赖性就会上升，因为阻止发布额外工作的是 exec 依赖性。这里的数字 16% 意味着大约 5/6 的时间，当出现停顿情况时，它不是由于 exec 依赖性造成的。换句话说，即使我们在这个内核中有很多执行依赖性，但大多数时候出现停顿，并不是因为 GPU 想转到下一行代码来发出——它是为了一些其他原因。
总结：
似乎至少有 2 个问题，都与不同类型的延迟有关：
- 在非常小的内核大小（例如 16384 个总线程）下，内核执行时间很短，因此测量会被例如内核启动延迟和可能的测量延迟。
- 内核大小非常小，不会使 GPU 尽可能多地进行并行工作，因此 IPC 和
sm_efficiency 之类的东西低于它们需要的水平，并且停滞原因: exec依赖比较高
任何时候你看到 sm_efficiency 这么低，一个可能的结论是没有足够的并行工作暴露给 GPU，因此计算吞吐量和内存都不是限制因素，而是延迟是性能的限制因素。
这与 the analysis-driven optimization logic（幻灯片 46 及以后）一致
并且可以通过向 GPU 公开更多工作来纠正。
我创建了三个合成 CUDA 内核，它们几乎都只进行算术运算。所有三个内核都是相同的，只是它们中的每一个都执行不同数量的操作。内核 #1 执行 8 次操作，内核 #2 执行 16 次操作，内核 #3 执行 32 次。以下是这三者的 CUDA 内核实现。
内核#1：
#ifndef kernelWGSXMAPIXLLXOPS8_H_
#define kernelWGSXMAPIXLLXOPS8_H_
__global__ void WGSXMAPIXLLXOPS8 (const float *GIn, float *GOut, const float M, const float N, const float P) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float MF = (float) M;
float NF = (float) N;
float PF = (float) P;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
... // 8 FMA operations
temp1 = temp1 * MF + temp1;
GOut[gid] = temp1;
}
}
void WGSXMAPIXLLXOPS8_wrapper (const float *GIn, float *GOut,
const float M, const float N, const float P,
int numBlocks, int threadPerBlock) {
WGSXMAPIXLLXOPS8<<<numBlocks, threadPerBlock>>> (GIn, GOut, M, N, P);
}
#endif
内核#2：
#ifndef kernelWGSXMAPIXLLXOPS16_H_
#define kernelWGSXMAPIXLLXOPS16_H_
__global__ void WGSXMAPIXLLXOPS16 (const float *GIn, float *GOut, const float M, const float N, const float P) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float MF = (float) M;
float NF = (float) N;
float PF = (float) P;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
... // 16 FMA operations
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
GOut[gid] = temp1;
}
}
void WGSXMAPIXLLXOPS16_wrapper (const float *GIn, float *GOut,
const float M, const float N, const float P,
int numBlocks, int threadPerBlock) {
WGSXMAPIXLLXOPS16<<<numBlocks, threadPerBlock>>> (GIn, GOut, M, N, P);
}
#endif
内核 #3：
#ifndef kernelWGSXMAPIXLLXOPS32_H_
#define kernelWGSXMAPIXLLXOPS32_H_
__global__ void WGSXMAPIXLLXOPS32 (const float *GIn, float *GOut, const float M, const float N, const float P) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float MF = (float) M;
float NF = (float) N;
float PF = (float) P;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
... // 32 FMA operations
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
GOut[gid] = temp1;
}
}
void WGSXMAPIXLLXOPS32_wrapper (const float *GIn, float *GOut,
const float M, const float N, const float P,
int numBlocks, int threadPerBlock) {
WGSXMAPIXLLXOPS32<<<numBlocks, threadPerBlock>>> (GIn, GOut, M, N, P);
}
#endif
线程总数已设置为 16384，块大小为 256。我计算了每个内核的总 GFlops，分别为 20.44、56.53 和 110.12 GFlops。我试图想出一个解释，但什么也想不起来。所以我尝试使用 nvprof 并监控所有指标。所有指标几乎相等，以下是一些对我来说很重要的指标（我还包括内核 1 到 3 的结果）：
sm_efficiency_instance: 14.99, 16.78, 19.82 %
ipc_instance: 0.57 , 0.93 , 1.53
inst_replay_overhead: 0.399, 0.268, 0.165
dram_write_throughput: 18.08, 17.72, 16.9 GB/s
issued_ipc: 0.99 , 1.18 , 1.52
issue_slot_utilization: 19.48, 24.64, 33.76 %
stall_exec_dependency: 21.84, 26.38, 42.95 %
很明显，它们都具有相同的 dram_write_throughput，因为它们都向 DRAM 写入相同数量的数据，并且线程总数也相同。我不明白的是sm_efficiency。我的内核都在做算术（一样），怎么他们的sm_efficiency不一样。另外，为什么在同一个内核中有更多的算法会提高效率？我的理解是，他们都应该有同样的问题来寻找 warp 以定位在 SM 上。
任何人都可以使用以下指标帮助我了解 GFlops 的区别吗？
基本问题是您没有 "saturated" 可用的 GPU。有各种与内核启动相关的开销。如果内核花费在计算上的时间与这个开销相比很小，那么你的计算就会被开销所扭曲。
T = 开销时间(OT) + 计算时间(CT)
Flops/s = Flops/T = 翻牌/(OT + CT)
如果计算时间与开销时间相比较小（您的内核就是这种情况），那么您的计算将受到开销时间的影响。另一方面，如果与开销相比计算时间足够大，那么开销对结果的影响相对较小。
这是一个完整的测试案例，有几个案例运行，CUDA 9.1，Tesla P100 PCIE：
$ cat t79.cu
#ifndef SLEN
#define SLEN (8)
#endif
#ifndef NTPB
#define NTPB (256)
#endif
#ifndef BLKS
#define BLKS (16384/NTPB)
#endif
const size_t blks = BLKS;
const size_t ntpb = NTPB;
typedef float Ftype;
#include <iostream>
template <int LEN>
__global__ void WGSXMAPIXLLXOPS (Ftype *GOut, const Ftype M) {
int gid = blockIdx.x * blockDim.x + threadIdx.x;
Ftype MF = (Ftype) M;
for (int lcdd = 0; lcdd < 1; lcdd++) {
float temp1 = 1.0;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
if (LEN > 8){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 16){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 32){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 64){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 128){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
if (LEN > 256){
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;
temp1 = temp1 * MF + temp1;}
#ifdef NO_WRITE
if (temp1 == -10.0)
#endif
GOut[gid] = temp1;
}
}
int main(){
float et;
Ftype *GOut;
const Ftype M = 1.0;
cudaMalloc(&GOut, blks*ntpb*sizeof(Ftype));
cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
WGSXMAPIXLLXOPS<SLEN><<<blks, ntpb>>> (GOut, M);
cudaDeviceSynchronize();
cudaEventRecord(start);
WGSXMAPIXLLXOPS<SLEN><<<blks, ntpb>>> (GOut, M);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&et, start, stop);
unsigned long long flpcnt = SLEN*2*blks*ntpb;
float Kflops_s = flpcnt/et;
std::cout << "MFlops per sec: " << Kflops_s/1000 << " kernel time: " << et << "ms" << std::endl;
cudaDeviceSynchronize();
}
$ nvcc -arch=sm_60 -o t79 t79.cu
$ ./t79
MFlops per sec: 14371.9 kernel time: 0.01824ms
$ nvprof ./t79
==14676== NVPROF is profiling process 14676, command: ./t79
MFlops per sec: 10101.1 kernel time: 0.025952ms
==14676== Profiling application: ./t79
==14676== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 100.00% 3.2320us 2 1.6160us 1.2480us 1.9840us void WGSXMAPIXLLXOPS<int=8>(float*, float)
API calls: 98.31% 389.62ms 1 389.62ms 389.62ms 389.62ms cudaMalloc
1.10% 4.3574ms 376 11.588us 357ns 465.31us cuDeviceGetAttribute
0.42% 1.6829ms 4 420.73us 272.19us 642.45us cuDeviceTotalMem
0.12% 487.27us 4 121.82us 90.094us 164.09us cuDeviceGetName
0.02% 80.363us 2 40.181us 15.789us 64.574us cudaLaunch
0.00% 17.118us 2 8.5590us 8.1400us 8.9780us cudaDeviceSynchronize
0.00% 13.118us 2 6.5590us 5.4290us 7.6890us cudaEventRecord
0.00% 10.603us 2 5.3010us 1.2440us 9.3590us cudaEventCreate
0.00% 8.5080us 8 1.0630us 460ns 1.7500us cuDeviceGet
0.00% 8.4590us 1 8.4590us 8.4590us 8.4590us cudaEventElapsedTime
0.00% 7.1350us 1 7.1350us 7.1350us 7.1350us cudaEventSynchronize
0.00% 6.8430us 4 1.7100us 180ns 5.9720us cudaSetupArgument
0.00% 4.7800us 3 1.5930us 437ns 2.8480us cuDeviceGetCount
0.00% 2.3490us 2 1.1740us 361ns 1.9880us cudaConfigureCall
$ nvcc -arch=sm_60 -o t79 t79.cu -DSLEN=512 -DBLKS=32768 -DNTPB=1024
$ ./t79
MFlops per sec: 8.08072e+06 kernel time: 4.25206ms
$
$ nvprof --metrics sm_efficiency_instance,ipc_instance,issued_ipc,issue_slot_utilization,stall_exec_dependency ./t79
==15447== NVPROF is profiling process 15447, command: ./t79
==15447== Some kernel(s) will be replayed on device 0 in order to collect all events/metrics.
Replaying kernel "void WGSXMAPIXLLXOPS<int=512>(float*, float)" (done)
Replaying kernel "void WGSXMAPIXLLXOPS<int=512>(float*, float)" (done)
MFlops per sec: 193432 kernel time: 177.632ms
==15447== Profiling application: ./t79
==15447== Profiling result:
==15447== Metric result:
Invocations Metric Name Metric Description Min Max Avg
Device "Tesla P100-PCIE-16GB (0)"
Kernel: void WGSXMAPIXLLXOPS<int=512>(float*, float)
2 issued_ipc Issued IPC 1.972106 1.972388 1.972247
2 issue_slot_utilization Issue Slot Utilization 98.23% 98.24% 98.24%
2 stall_exec_dependency Issue Stall Reasons (Execution Dependency) 16.35% 16.36% 16.36%
2 ipc Executed IPC 1.971976 1.972254 1.972115
2 sm_efficiency Multiprocessor Activity 99.78% 99.78% 99.78%
$
第一个 运行，数字与您的匹配（16384 个线程，每个块 256 个线程，8 条 FFMA 指令）显示内核持续时间约为 17us。然而，当我们在分析器中 运行 这种情况时，我们观察到实际内核执行只有大约 1.5us，剩下的是各种开销，包括内核启动延迟，以及使用 cudaEvent系统进行计时。所以这会使数字偏离。
另一方面，当我们启动大量块、每个块的线程数以及每个线程工作时，我们得到的数字是 P100 峰值能力的 80%。
当您从内核 1 到内核 3 时，您的大多数指标都在增加（变得更好）（dram 吞吐量除外，这是合理的。随着内核时间的增加，对于相同数量的写入数据，dram 平均吞吐量下跌降落）。这与为 GPU 提供更多工作是一致的，因此它可以隐藏各种延迟并分摊大量工作的开销。
让我们看一下上述最终 运行/"large" 内核的一些指标：
2 issued_ipc Issued IPC 1.972106 1.972388 1.972247
2 issue_slot_utilization Issue Slot Utilization 98.23% 98.24% 98.24%
2 stall_exec_dependency Issue Stall Reasons (Execution Dependency) 16.35% 16.36% 16.36%
2 ipc Executed IPC 1.971976 1.972254 1.972115
2 sm_efficiency Multiprocessor Activity 99.78% 99.78% 99.78%
IPC 大约为每个时钟 2，高于您的内核 3。请注意，此处的 IPC 为 2 是一个合理的上限：sm_60 SM 有 64 single-precision 个单元，足以安排每个时钟 2 条 FFMA 指令。
SM 效率和 issue_slot_utilization 是相似的指标。这意味着大约 98% 的时间，SM 可以在任何给定的时钟周期内发出一条或多条指令。
stall（exec 依赖项）正在回答问题，"across all the actual stall situations, what percent were due to execution dependency?"。您的内核在每行源代码之间具有执行依赖性——因为每行都依赖于前一行的结果。这意味着在汇编级别，每条 FFMA 指令都将取决于前一条指令的结果，因此只有在前一条指令完成后才能发出。
如果 SM 的可用工作订阅不足，那么 stall exec 依赖性就会上升，因为阻止发布额外工作的是 exec 依赖性。这里的数字 16% 意味着大约 5/6 的时间，当出现停顿情况时，它不是由于 exec 依赖性造成的。换句话说，即使我们在这个内核中有很多执行依赖性，但大多数时候出现停顿，并不是因为 GPU 想转到下一行代码来发出——它是为了一些其他原因。
总结：
似乎至少有 2 个问题，都与不同类型的延迟有关：
- 在非常小的内核大小（例如 16384 个总线程）下，内核执行时间很短，因此测量会被例如内核启动延迟和可能的测量延迟。
- 内核大小非常小，不会使 GPU 尽可能多地进行并行工作，因此 IPC 和
sm_efficiency之类的东西低于它们需要的水平，并且停滞原因: exec依赖比较高
任何时候你看到 sm_efficiency 这么低，一个可能的结论是没有足够的并行工作暴露给 GPU，因此计算吞吐量和内存都不是限制因素，而是延迟是性能的限制因素。
这与 the analysis-driven optimization logic（幻灯片 46 及以后）一致
并且可以通过向 GPU 公开更多工作来纠正。