0

我使用 nvprof 在 P100 上分析了一个简单的 vecadd 示例（n=1024），但观察到 dram_write_bytes 只有 256（而不是我预期的 1024*4）。有人可以解释为什么这个数字很小吗？我需要添加哪些其他指标来计算全局内存写入？谢谢。float_count_sp 数字是正确的 (1024)。

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

__global__ void vecAdd(float* a, float* b, float* c, int n){
    int id = blockIdx.x*blockDim.x + threadIdx.x;
    if(id < n) c[id] = a[id] + b[id];
}

int main(int argc, char* argv[]){
    int n = 1024;
    float *h_a, *d_a;
    float *h_b, *d_b;
    float *h_c, *d_c;
    size_t bytes = n*sizeof(float);
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    int i;
    for(i = 0; i < n; i++){
        h_a[i] = sin(i)*sin(i);
        h_b[i] = cos(i)*cos(i+1);
    }
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    vecAdd <<<1, 1024>>> (d_a, d_b, d_c, n);
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    float sum = 0;
    for(i = 0; i < n; i++)
        sum += h_c[i] - h_a[i] - h_b[i];
    printf("final diff: %f\n", sum/n);
    
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

是不是和nvprof的采样有关？有一次我得到 384 个字节。有时我什至得到 0 个字节。奇怪的是：如果我将 n 更改为 1024*1024，我得到的字节数比我预期的要多（4688032）。4688032/1024/1024/4 = 1.11。

4

1 回答 1

3

没有遵守您的期望并且数据正在发生变化的原因有几个：

  1. GPU 内存系统由所有引擎共享。主要引擎是图形/计算引擎，但其他引擎（如复制引擎、显示等）访问设备内存和内存控制（FB = 帧缓冲区）计数器没有跟踪请求者的方法。

  2. NVPROF 注入不会尝试从 L2 缓存中逐出所有上下文内存。启动前的 cudaMemcpys 和 nvprof 中的内核重放代码将使 L2 缓存处于不一致状态。

  3. 4KB 的初始大小对于准确跟踪来说太小了。完整的数据集可以来自 cudaMemcpy 或重放的 L2。此外，您看到的字节可能来自其他客户端，例如常量缓存。

强烈建议您将缓冲区大小调整到合理的大小。在较新的 GPU 上，Nsight Compute 分析器改进了各种客户端的 L2 级别细分，以帮助检测意外流量。此外，Nsight Compute 重放逻辑会清除 L2 缓存，以便每次重放都有一致的开始状态。

如果您连接了显示器，建议在查看 DRAM 计数器时将显示器移至不同的 GPU。nvprof L2 计数器通常按来自 SM 的流量过滤计数，因此来自复制引擎、显示控制器、MMU、常量缓存等的流量不会显示在 L2 计数器中。

于 2020-07-14T18:29:48.970 回答