CUDA中的常量内存__constant__和cudaMemcpyToSymbol-CSDN博客

__constant__声明内存为常量内存

使用常量内存可以提升运算性能的原因如下：

对常量内存的单次读操作可以广播到其他的“邻近(nearby)”线程，这将节约15次读取操作；
高速缓存。常量内存的数据将缓存起来，因此对于相同地址的连续操作将不会产生额外的内存通信量；

在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。
当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在半线程束中包含16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量数据，那么这种方式产生的内存流量只是使用全局内存时的1/16。

常量内存的拷贝必须用cudaMemcpyToSymbol，如下

cudaGetSymbolAddress()用来获得__constant__内存的全局物理地址；

cudaGetSymbolSize()可以获得__constant__内存大小

    float *dp = NULL;
    cudaGetSymbolAddress((void**)&dp,devData);
    cudaMemcpy(dp,&value,sizeof(float),cudaMemcpyHostToDevice);
    // invoke the kernel
    checkGlobalVariable<<<1, 1>>>();
    // copy the global variable back to the host
    CHECK(cudaMemcpyFromSymbol(&value, devData, sizeof(float)));
    printf("Host:   the value changed by the kernel to %f\n", value);