CUDA 不能使用所有可用的常量内存
CUDA can't use all available constant memory
我有一个使用协作组执行某些操作的代码。因此我编译我的代码：
/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 = ¶m;
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 字节的常量内存，原因不明。
这个答案是推测性的，因为 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 error : File uses too much global constant data (0x10018 bytes, 0x10000 max)
如果我们改用双精度函数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 error : File uses too much global constant data (0x10110 bytes, 0x10000 max)
为了记录这个用例中到底发生了什么，我拼凑了编译过程中的以下工作。希望它能阐明这个问题是如何产生的，以及一些有用的诊断工具，同时消除一些误解。
请注意，这是一项正在进行的工作，随着更多信息的出现，可能会定期更新。请根据您的需要进行编辑和投稿
首先，如评论中所述，完全有可能分配常量内存的每个字节直到 64kb 限制。这个例子几乎就是原始问题中描述的用例：
const int sz = 65536;
const int NMax = sz / sizeof(float);
__constant__ float buffer[NMax];
__global__
void akernel(const float* __restrict__ arg1, float* __restrict__ arg2, int N)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) {
float ans = 0;
#pragma unroll 128
for(int i=0; i<NMax; i++) {
float val = buffer[i];
float y = (i%2 == 0) ? 1.f : -1.f;
float x = val / 255.f;
ans = ans + y * sinf(x);
}
arg2[tid] = ans + arg1[tid];
}
}
并且编译没有问题（Godbolt link here）。这证明问题中的 linker 阶段必须从其他代码中引入额外的常量内存分配，无论是用户代码、其他设备库还是设备 运行time 支持。
因此，让我们将注意力转向更新问题中 posted 的重现案例，稍微修改一下，以便通过稍微减少常量内存占用量来通过编译和 link 阶段， 64536 字节的缓冲区：
$ nvcc -arch=sm_61 --device-c -g -O2 -Xptxas="-v" -o constmemuse.cu.o constmemuse.cu
constmemuse.cu(51): warning: variable "centers" was declared but never referenced
constmemuse.cu(52): warning: variable "centersDist" was declared but never referenced
constmemuse.cu(31): warning: variable "dist" was set but never used
detected during instantiation of "void getNeatestCenter<R,C,bRO,ROWMAJ>(foo_params) [with R=128, C=64, bRO=true, ROWMAJ=true]"
constmemuse.cu(26): warning: variable "mindist" was declared but never referenced
detected during instantiation of "void getNeatestCenter<R,C,bRO,ROWMAJ>(foo_params) [with R=128, C=64, bRO=true, ROWMAJ=true]"
ptxas info : 0 bytes gmem, 64536 bytes cmem[3]
ptxas info : Function properties for cudaDeviceGetAttribute
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Compiling entry function '_Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params' for 'sm_61'
ptxas info : Function properties for _Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 5 registers, 360 bytes cmem[0]
ptxas info : Function properties for cudaMalloc
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessor
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaGetDevice
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Compiling entry function '_Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params' for 'sm_61'
ptxas info : Function properties for _Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 5 registers, 360 bytes cmem[0]
ptxas info : Function properties for cudaFuncGetAttributes
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
几点：
64536 bytes cmem[3] 显示用户可控常量存储区的大小，正如我们指定的那样 ptxas info : Used 5 registers, 360 bytes cmem[0] 显示函数的寄存器用法，cmem[0] 是内部保留的常量内存库，用于保存内核参数和编译器放入常量内存的任何其他内容。请注意，寄存器溢出会进入本地内存，而不是常量内存。
那么现在让我们 运行 设备 linking 步骤：
$ nvcc -arch=sm_61 -gencode=arch=compute_61,code=sm_61, -g -dlink -Xnvlink="-v" -o constmemuse.o constmemuse.cu.o
nvlink info : 9944 bytes gmem, 64792 bytes cmem[3] (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 20 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 23 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 28 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 23 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 12 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 17 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 14 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 16 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 16 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 14 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 17 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 8 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 11 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 12 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 11 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 21 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '_Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 360 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '_Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 360 bytes cmem[0], 0 bytes lmem (target: sm_61)
一些补充说明：
9944 bytes gmem, 64792 bytes cmem[3] 现在显示 linked 模块的全局和常量内存预留。如您所见，我们在常量组 0（用户可修改组）中继承了 256 个额外字节，以及 9944 个字节的静态保留全局内存。如果数组分配为 65536 字节，如问题所示，linkage 将失败，因为它超过了 64kb 的限制。- 您可以看到许多设备 运行time 库函数已在 linkage 阶段（memcpy 和 memset）link自动编辑
很明显，额外的常量内存使用即将到来 link 设备 运行 时间，可以通过 cuobjdump post hoc 确认。编译对象：
$ cuobjdump -res-usage constmemuse.cu.o
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
Resource usage:
Common:
GLOBAL:0 CONSTANT[3]:64536
Function cudaDeviceGetAttribute:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function _Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params:
REG:5 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaMalloc:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaOccupancyMaxActiveBlocksPerMultiprocessor:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaGetDevice:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function _Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params:
REG:5 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaFuncGetAttributes:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Fatbin ptx code:
================
arch = sm_61
code version = [6,4]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
ptxasOptions = -v --compile-only
和 linking 之后的对象：
$ cuobjdump -res-usage constmemuse.o
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
Resource usage:
Common:
GLOBAL:9944 CONSTANT[3]:64792
Function _Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function _Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:21 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:11 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:12 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:11 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:8 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:17 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:14 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:16 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:16 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:14 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:17 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:12 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:23 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:28 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:23 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:20 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaCGGetIntrinsicHandle:
REG:6 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
在接受的答案中已经证明，数学库可以为某些三角函数和超越函数的系数和查找表保留常量内存。然而，在这种情况下，原因似乎是内核中使用协作组发出的支持样板。进一步研究额外的 bank 0 常量内存的确切来源需要对该代码进行反汇编和逆向工程，我现在不打算这样做。
我有一个使用协作组执行某些操作的代码。因此我编译我的代码：
/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 = ¶m;
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 字节的常量内存，原因不明。
这个答案是推测性的，因为 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 error : File uses too much global constant data (0x10018 bytes, 0x10000 max)
如果我们改用双精度函数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 error : File uses too much global constant data (0x10110 bytes, 0x10000 max)
为了记录这个用例中到底发生了什么，我拼凑了编译过程中的以下工作。希望它能阐明这个问题是如何产生的，以及一些有用的诊断工具，同时消除一些误解。
请注意，这是一项正在进行的工作，随着更多信息的出现，可能会定期更新。请根据您的需要进行编辑和投稿
首先，如评论中所述，完全有可能分配常量内存的每个字节直到 64kb 限制。这个例子几乎就是原始问题中描述的用例：
const int sz = 65536;
const int NMax = sz / sizeof(float);
__constant__ float buffer[NMax];
__global__
void akernel(const float* __restrict__ arg1, float* __restrict__ arg2, int N)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) {
float ans = 0;
#pragma unroll 128
for(int i=0; i<NMax; i++) {
float val = buffer[i];
float y = (i%2 == 0) ? 1.f : -1.f;
float x = val / 255.f;
ans = ans + y * sinf(x);
}
arg2[tid] = ans + arg1[tid];
}
}
并且编译没有问题（Godbolt link here）。这证明问题中的 linker 阶段必须从其他代码中引入额外的常量内存分配，无论是用户代码、其他设备库还是设备 运行time 支持。
因此，让我们将注意力转向更新问题中 posted 的重现案例，稍微修改一下，以便通过稍微减少常量内存占用量来通过编译和 link 阶段， 64536 字节的缓冲区：
$ nvcc -arch=sm_61 --device-c -g -O2 -Xptxas="-v" -o constmemuse.cu.o constmemuse.cu
constmemuse.cu(51): warning: variable "centers" was declared but never referenced
constmemuse.cu(52): warning: variable "centersDist" was declared but never referenced
constmemuse.cu(31): warning: variable "dist" was set but never used
detected during instantiation of "void getNeatestCenter<R,C,bRO,ROWMAJ>(foo_params) [with R=128, C=64, bRO=true, ROWMAJ=true]"
constmemuse.cu(26): warning: variable "mindist" was declared but never referenced
detected during instantiation of "void getNeatestCenter<R,C,bRO,ROWMAJ>(foo_params) [with R=128, C=64, bRO=true, ROWMAJ=true]"
ptxas info : 0 bytes gmem, 64536 bytes cmem[3]
ptxas info : Function properties for cudaDeviceGetAttribute
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Compiling entry function '_Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params' for 'sm_61'
ptxas info : Function properties for _Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 5 registers, 360 bytes cmem[0]
ptxas info : Function properties for cudaMalloc
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessor
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaGetDevice
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Compiling entry function '_Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params' for 'sm_61'
ptxas info : Function properties for _Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 5 registers, 360 bytes cmem[0]
ptxas info : Function properties for cudaFuncGetAttributes
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
几点：
64536 bytes cmem[3]显示用户可控常量存储区的大小，正如我们指定的那样ptxas info : Used 5 registers, 360 bytes cmem[0]显示函数的寄存器用法，cmem[0]是内部保留的常量内存库，用于保存内核参数和编译器放入常量内存的任何其他内容。请注意，寄存器溢出会进入本地内存，而不是常量内存。
那么现在让我们 运行 设备 linking 步骤：
$ nvcc -arch=sm_61 -gencode=arch=compute_61,code=sm_61, -g -dlink -Xnvlink="-v" -o constmemuse.o constmemuse.cu.o
nvlink info : 9944 bytes gmem, 64792 bytes cmem[3] (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 20 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 23 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 28 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 23 registers, 0 stack, 2056 bytes smem, 448 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 10 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 12 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 17 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 14 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_': (target: sm_61)
nvlink info : used 16 registers, 0 stack, 2056 bytes smem, 416 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 16 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 14 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 17 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 8 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 11 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 12 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 11 registers, 0 stack, 0 bytes smem, 400 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '__nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_': (target: sm_61)
nvlink info : used 21 registers, 0 stack, 0 bytes smem, 424 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '_Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 360 bytes cmem[0], 0 bytes lmem (target: sm_61)
nvlink info : Function properties for '_Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params': (target: sm_61)
nvlink info : used 6 registers, 0 stack, 0 bytes smem, 360 bytes cmem[0], 0 bytes lmem (target: sm_61)
一些补充说明：
9944 bytes gmem, 64792 bytes cmem[3]现在显示 linked 模块的全局和常量内存预留。如您所见，我们在常量组 0（用户可修改组）中继承了 256 个额外字节，以及 9944 个字节的静态保留全局内存。如果数组分配为 65536 字节，如问题所示，linkage 将失败，因为它超过了 64kb 的限制。- 您可以看到许多设备 运行time 库函数已在 linkage 阶段（memcpy 和 memset）link自动编辑
很明显，额外的常量内存使用即将到来 link 设备 运行 时间，可以通过 cuobjdump post hoc 确认。编译对象：
$ cuobjdump -res-usage constmemuse.cu.o
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
Resource usage:
Common:
GLOBAL:0 CONSTANT[3]:64536
Function cudaDeviceGetAttribute:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function _Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params:
REG:5 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaMalloc:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaOccupancyMaxActiveBlocksPerMultiprocessor:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaGetDevice:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function _Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params:
REG:5 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaFuncGetAttributes:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags:
REG:5 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
Fatbin ptx code:
================
arch = sm_61
code version = [6,4]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
ptxasOptions = -v --compile-only
和 linking 之后的对象：
$ cuobjdump -res-usage constmemuse.o
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
Resource usage:
Common:
GLOBAL:9944 CONSTANT[3]:64792
Function _Z16getNeatestCenterILi128ELi64ELb1ELb1EEv10foo_params:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function _Z16getNeatestCenterILi32ELi32ELb1ELb1EEv10foo_params:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:360 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:21 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:11 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:12 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:11 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:8 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceIjLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:400 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:17 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:14 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi0ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:16 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi0ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi0EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memset_3d_deviceImLi1ELi1ELi1EEvPhhjT_S1_S1_S1_S1_jjjjjjjS1_S0_:
REG:6 STACK:0 SHARED:0 LOCAL:0 CONSTANT[0]:424 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:16 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:14 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 CONSTANT[2]:4 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:17 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:12 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceIjLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:416 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:23 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:28 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:23 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi0ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:20 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi0ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi0EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function __nv_static_51__38_cuda_device_runtime_compute_75_cpp1_ii_8b1a5d37__Z16memcpy_3d_deviceImLi1ELi1ELi1EEvPKhPhT_S3_S3_S3_S3_S3_S3_jjjjjjjjS3_S1_S2_:
REG:10 STACK:0 SHARED:2056 LOCAL:0 CONSTANT[0]:448 TEXTURE:0 SURFACE:0 SAMPLER:0
Function cudaCGGetIntrinsicHandle:
REG:6 STACK:0 SHARED:0 LOCAL:0 TEXTURE:0 SURFACE:0 SAMPLER:0
在接受的答案中已经证明，数学库可以为某些三角函数和超越函数的系数和查找表保留常量内存。然而，在这种情况下，原因似乎是内核中使用协作组发出的支持样板。进一步研究额外的 bank 0 常量内存的确切来源需要对该代码进行反汇编和逆向工程，我现在不打算这样做。