cuda - 给GPU带来负担的小程序？

我在这里草拟一个可能的解决方案。您将需要一些实验来最大化 GPU 的热负载。一般来说，数据移动在能量上是昂贵的，比现代处理器中的计算要昂贵得多。因此，对大量数据进行洗牌会增加功耗。同时，我们希望计算单元对功耗的贡献增加。乘数往往是最大的耗电者；在现代处理器中，我们可能希望以 FMA（融合乘加）单元为目标。

各种 GPU 的双精度数学运算吞吐量较低，其他 GPU 的半精度数学运算吞吐量较低。因此，我们希望专注于负载计算部分的单精度数学。我们希望能够轻松更改计算与内存活动的比率。一种方法是使用以霍纳方案为基本构建块的多项式的展开评估，使用POLY_DEPTH步骤。REPS我们在一个循环中重复时间。在循环之前，我们从全局内存中检索源数据，在循环终止后，我们将结果存储到全局内存中。通过改变REPS我们可以尝试不同的计算/内存平衡设置。

人们可以进一步试验指令级并行性、数据模式（因为乘法器的功耗通常因位模式而异），并通过使用 CUDA 流来实现内核执行和 PCIe 数据传输的重叠来添加 PCIe 活动。下面我只是使用了一些随机常数作为乘数数据。

显然，我们希望用大量线程填充 GPU。为此，我们可以使用一个相当小的THREADS_PER_BLK值，为我们填充每个 SM 提供细粒度。我们可能希望将块数选择为 SM 数量的倍数以尽可能均匀地分配负载，或者使用一个MAX_BLOCKS平均除以常见 SM 计数的值。我们应该接触多少源内存和目标内存取决于实验：我们可以将LEN元素数组定义为块数的倍数。最后，我们要执行这样定义和配置的内核ITER次数，以创建一段时间的连续负载。

请注意，当我们施加负载时，GPU 会发热，这反过来会进一步增加其功耗。要实现最大热负载，需要运行负载生成应用程序 5 分钟或更长时间。进一步注意，GPU 电源管理可能会动态降低时钟频率和电压以降低功耗，并且功率上限可能会在您达到热限制之前启动。根据 GPU，您可以将功率上限设置为高于该nvidia-smi实用程序默认使用的功率上限。

根据 TechPowerUp 的 GPU-Z 实用程序的报告，下面的程序使我的 Quadro P2000 保持在功率上限，GPU 负载为 98%，内存控制器负载为 83%-86%。它肯定需要对其他 GPU 进行调整。

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

#define THREADS_PER_BLK (128)
#define MAX_BLOCKS      (65520)
#define LEN             (MAX_BLOCKS * 1024)
#define POLY_DEPTH      (30)
#define REPS            (2)
#define ITER            (100000)

// Macro to catch CUDA errors in CUDA runtime calls
#define CUDA_SAFE_CALL(call)                                          \
do {                                                                  \
    cudaError_t err = call;                                           \
    if (cudaSuccess != err) {                                         \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString(err) );       \
        exit(EXIT_FAILURE);                                           \
    }                                                                 \
} while (0)

// Macro to catch CUDA errors in kernel launches
#define CHECK_LAUNCH_ERROR()                                          \
do {                                                                  \
    /* Check synchronous errors, i.e. pre-launch */                   \
    cudaError_t err = cudaGetLastError();                             \
    if (cudaSuccess != err) {                                         \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString(err) );       \
        exit(EXIT_FAILURE);                                           \
    }                                                                 \
    /* Check asynchronous errors, i.e. kernel failed (ULF) */         \
    err = cudaDeviceSynchronize();                                    \
    if (cudaSuccess != err) {                                         \
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.\n",\
                 __FILE__, __LINE__, cudaGetErrorString( err) );      \
        exit(EXIT_FAILURE);                                           \
    }                                                                 \
} while (0)

__global__ void burn (const float * __restrict__ src, 
                      float * __restrict__ dst, int len)
{
    int stride = gridDim.x * blockDim.x;
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    for (int i = tid; i < len; i += stride) {
        float p = src[i] + 1.0;
        float q = src[i] + 3.0f;
        for (int k = 0; k < REPS; k++) {
#pragma unroll POLY_DEPTH
            for (int j = 0; j < POLY_DEPTH; j++) {
                p = fmaf (p, 0.68073987f, 0.8947237f);
                q = fmaf (q, 0.54639739f, 0.9587058f);
            }
        }
        dst[i] = p + q;
    }
}    

int main (int argc, char *argv[])
{
    float *d_a, *d_b;

    /* Allocate memory on device */
    CUDA_SAFE_CALL (cudaMalloc((void**)&d_a, sizeof(d_a[0]) * LEN));
    CUDA_SAFE_CALL (cudaMalloc((void**)&d_b, sizeof(d_b[0]) * LEN));

    /* Initialize device memory */
    CUDA_SAFE_CALL (cudaMemset(d_a, 0x00, sizeof(d_a[0]) * LEN)); // zero
    CUDA_SAFE_CALL (cudaMemset(d_b, 0xff, sizeof(d_b[0]) * LEN)); // NaN

    /* Compute execution configuration */
    dim3 dimBlock(THREADS_PER_BLK);
    int threadBlocks = (LEN + (dimBlock.x - 1)) / dimBlock.x;
    if (threadBlocks > MAX_BLOCKS) threadBlocks = MAX_BLOCKS;
    dim3 dimGrid(threadBlocks);

    printf ("burn: using %d threads per block, %d blocks, %f GB\n", 
            dimBlock.x, dimGrid.x, 2e-9*LEN*sizeof(d_a[0]));

    for (int k = 0; k < ITER; k++) {
        burn<<<dimGrid,dimBlock>>>(d_a, d_b, LEN);
        CHECK_LAUNCH_ERROR();
    }

    CUDA_SAFE_CALL (cudaFree(d_a));
    CUDA_SAFE_CALL (cudaFree(d_b));

    return EXIT_SUCCESS;
}

cuda - 给GPU带来负担的小程序？

问题描述

解决方案

推荐阅读