CUDA 不能使用所有可用的常量内存

Question

我有一段代码使用协作组来执行一些操作。因此我编译我的代码:

/usr/local/cuda/bin/nvcc -arch=sm_61 -gencode=arch=compute_61,code=sm_61, --device-c -g -O2 foo.cu

然后我尝试调用设备链接器:

/usr/local/cuda/bin/nvcc -arch=sm_61 -gencode=arch=compute_61,code=sm_61, -g -dlink foo.o

然后它会产生错误:

ptxas error : File uses too much global constant data (0x10100 bytes, 0x10000 max)

问题是由我分配常量内存的方式引起的:

__constant__ float d_cnst_centers[CONST_MEM / sizeof(float)];

其中 CONST_MEM = 65536 字节，这是我从 SM_61 的设备查询中获得的。但是，如果我将常量内存减少到 64536 之类的东西，问题就消失了。几乎就像在编译期间出于某些目的“保留”常量内存一样。我搜索了 CUDA 文档，但没有找到令人满意的答案。使用可用的最大常量内存是否安全？为什么会出现这个问题？

编辑:这是触发 SM_61 错误的代码片段:

#include <algorithm>
#include <vector>
#include <type_traits>
#include <cuda_runtime.h>
#include <cfloat>
#include <iostream>

#include <cooperative_groups.h>

using namespace cooperative_groups;


struct foo_params {
    float * points;
    float * centers;
    int * centersDist;
    int * centersIndex;
    int numPoints;
};

__constant__ float d_cnst_centers[65536 / sizeof(float)];

template <int R, int C>
__device__ int 
nearestCenter(float * points, float * pC) {
    float mindist = FLT_MAX;
    int minidx = 0;
    int clistidx = 0;
    for(int i=0; i<C;i++) {
        clistidx = i*R;
        float dist;
        {
            float *point = points;
            float *center = &pC[clistidx];
            float accum;
            for(int i = 0; i<R; i++) {
                float delta = point[i] - center[i];
                accum += delta*delta;
            }
            dist = sqrt(accum);
        }
        /* ... */
    }
    return minidx;
}


template<int R, int C, bool bRO, bool ROWMAJ=true>
__global__ void getNeatestCenter(struct foo_params params) {
        float * points = params.points;
        float * centers = params.centers;
        int * centersDist = params.centersDist;
        int * centersIndex = params.centersIndex;
        int numPoints = params.numPoints;

        grid_group grid = this_grid();
        {
            int idx = blockIdx.x*blockDim.x+threadIdx.x;
            if (idx < numPoints) {
                centersIndex[idx] = nearestCenter<R,C>(&points[idx*R], d_cnst_centers);
            }
        }
        /* ... other code */
}

int main () {
    // foo paramaters, for illustration purposes
    struct foo_params param;
    param.points = NULL;
    param.centers = NULL;
    param.centersDist = NULL;
    param.centersIndex = NULL;
    param.numPoints = 1000000;
    void *p_params = &param;

    int minGridSize = 0, blockSize = 0;
    cudaOccupancyMaxPotentialBlockSize(
                            &minGridSize,
                            &blockSize,
                            (void*)getNeatestCenter<128, 64, true>,
                            0,
                            0);

    dim3 dimGrid(minGridSize, 1, 1), dimBlock(blockSize, 1, 1);

    cudaLaunchCooperativeKernel((void *)getNeatestCenter<32, 32, true>, dimGrid, dimBlock, &p_params);
}

问题似乎是由线路引起的:

grid_group grid = this_grid();

这似乎在没有已知原因的情况下使用了大约 0x100 字节的常量内存。

Answer 1

这个答案是推测性的，因为 OP 没有提供最少但完整的重现代码。

GPU 包含多个常量内存库，用于程序存储的不同部分。其中一个银行供程序员使用。重要的是，CUDA 标准数学库代码使用相同的库，因为数学库代码通过函数内联成为程序员代码的一部分。在过去，这是显而易见的，因为整个 CUDA 数学库最初只是几个头文件。

一些数学函数在内部需要小的常量数据表。具体示例是 sin、cos、tan。当使用这些数学函数时，程序员可用的 __constant__ 数据量从 64KB 减少了少量。以下是一些用于演示目的的示例程序，使用 CUDA 8 工具链和 -arch=sm_61 编译:

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

#define CONST_MEM (65536)
__constant__ float d_cnst_centers[CONST_MEM / sizeof(float)] = {1};

__global__ void kernel (int i, float f)
{
    float r = d_cnst_centers[i] * expf(f);
    printf ("r=%15.8f\n", r);
}

int main (void)
{
    kernel<<<1,1>>>(0,25.0f);
    cudaDeviceSynchronize();
    return EXIT_SUCCESS;
}

这可以很好地编译并在运行时打印 r=72004902912.00000000。现在让我们将 expf 更改为 sinf:

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

#define CONST_MEM (65536)
__constant__ float d_cnst_centers[CONST_MEM / sizeof(float)] = {1};

__global__ void kernel (int i, float f)
{
    float r = d_cnst_centers[i] * sinf(f);
    printf ("r=%15.8f\n", r);
}

int main (void)
{
    kernel<<<1,1>>>(0,25.0f);
    cudaDeviceSynchronize();
    return EXIT_SUCCESS;
}

这会在编译期间引发错误:ptxas 错误:文件使用了过多的全局常量数据(0x10018 字节，最大 0x10000)

如果我们改用 double 函数 sin，则需要更多常量内存:

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

#define CONST_MEM (65536)
__constant__ float d_cnst_centers[CONST_MEM / sizeof(float)] = {1};

__global__ void kernel (int i, float f)
{
    float r = d_cnst_centers[i] * sin((double)f);
    printf ("r=%15.8f\n", r);
}

int main (void)
{
    kernel<<<1,1>>>(0,25.0f);
    cudaDeviceSynchronize();
    return EXIT_SUCCESS;
}

我们收到错误信息:ptxas 错误:文件使用了过多的全局常量数据(0x10110 字节，最大 0x10000)