CUDA——常量内存初始化

介绍

常量内存有多种初始化方式：

__constant__ int arrA[10];
__constant__ int arrB[3] = {1,2,3};
__constant__ int C = 5;

int main(int argc, char** argv)
{
	......
	int h_arr[10] = {1,2,3,4,5,6,7,8,9,0};
	cudaMemcpyToSymbol(arrA, h_arr, sizeof(int) * 10);
}

但是有时候我们需要将一个变量声明为常量，并且这个值需要在执行host代码之后再确定的，例如：

__constant__ int D;
int main(int argc, char** argv)
{
	......
	int h_D;
	cin >> h_D;
	cudaMemcpyToSymbol(&D, &h_D, sizeof(int));
}

运行该代码会报错：

Error: test.cu:29, code: 13, reason: invalid device symbol

分析

一开始我以为是设备常量变量无法在主机代码中通过&符号直接解析出地址，于是我调用cudaGetSymbolAddress函数来获取常量D的地址：

int* tmp_ptr;
cudaGetSymbolAddress((void**)&tmp_ptr, D);
cudaMemcpyToSymbol(tmp_ptr, &tmp, sizeof(int))

结果依然报错。

答案

对变量进行常量声明之后，需要在主机代码中为其分配内存空间。

__constant__ int D;
int main(int argc, char** argv)
{
	......
	int h_D;
	cin >> h_D;
	cudaMalloc((void**)&D, sizeof(int));
	cudaMemcpyToSymbol(D, &h_D, sizeof(int));
}