0

我是一个新手 CUDA 程序员。我最近了解了更多关于在较低占用率下实现更好性能的信息。这是一个代码片段，我需要帮助来理解一些关于重播开销和指令级并行性的事情

__global__ void myKernel(double *d_dst, double *d_a1, double *d_a2, size_t SIZE) 
{

    int tId = threadIdx.x + blockDim.x * blockIdx.x;

    d_dst[tId]            = d_a1[tId] * d_a2[tId];
    d_dst[tId + SIZE]     = d_a1[tId + SIZE] * d_a2[tId + SIZE];
    d_dst[tId + SIZE * 2] = d_a1[tId + SIZE * 2] * d_a2[tId + SIZE * 2];
    d_dst[tId + SIZE * 3] = d_a1[tId + SIZE * 3] * d_a2[tId + SIZE * 3];
}

这是我的简单内核，它简单地将两个二维数组相乘以形成第三个二维数组（从逻辑角度来看），其中这些数组都作为平面一维数组放置在设备内存中。

下面我展示另一段代码片段：

void doCompute() {

    double *h_a1;
    double *h_a2;

    size_t SIZE = pow(31, 3) + 1;

   // Imagine h_a1, h_a2 as 2D arrays
   // with 4 rows and SIZE Columns
   // For convenience created as 1D arrays 

    h_a1 = (double *) malloc(SIZE * 4 * sizeof(double));
    h_a2 = (double *) malloc(SIZE * 4 * sizeof(double));

    memset(h_a1, 5.0, SIZE * 4 * sizeof(double));
    memset(h_a2, 5.0, SIZE * 4 * sizeof(double));

    double *d_dst;
    double *d_a1;
    double *d_a2;

    cudaMalloc(&d_dst, SIZE * 4 * sizeof(double));
    cudaMalloc(&d_a1,  SIZE * 4 * sizeof(double));
    cudaMalloc(&d_a2,  SIZE * 4 * sizeof(double));

    cudaMemcpy(d_a1, h_a1, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_a2, h_a2, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice);

    int BLOC_SIZE = 32;
    int GRID_SIZE = (SIZE + BLOC_SIZE - 1) / BLOC_SIZE;

    myKernel <<< GRID_SIZE, BLOC_SIZE >>> (d_dst, d_a1, d_a2, SIZE);
}

Q1）我在这里打破任何合并的内存访问模式吗？

Q2）我可以说对内存的访问，它们在内核中的编码方式也是指令级并行的例子吗？如果是，我使用的是 ILP2 还是 ILP4？为什么？

Q3) 如果我所做的一切都是正确的，那么为什么 nvvp 分析器会给我以下消息？

Total Replay Overhead: 4.6%
Global Cache Replay Overhead: 30.3%

如何减少或修复它们？

干杯，

4

1 回答 1

1

编译器为可能的 ILP 漏洞调度指令的能力有限。GPU 本身也必须具有 ILP 能力，其程度因 GPU 代而异。是的，任何不可用的资源都可能导致扭曲停止，典型的资源是内存中需要的数据。您所询问的重放量的定义在此处给出。

因此，例如，全局缓存重放开销将由缓存未命中触发，并且您的代码将有一些缓存未命中。即使您有 100% 的合并访问和（几乎）100% 的带宽利用效率，缓存未命中也是可能的。

于 2013-07-11T19:25:50.077 回答