同一常量内存数组在不同源文件上的使用

Question

我有一个__constant__内存阵列，其中包含许多内核所需的信息，这些内核放在不同的源文件中。这个常量内存数组是在头文件GlobalParameters.h中定义的，所有包含需要访问这个数组的内核的文件都会对它进行#included。

I just discovered (查看talonmies的答案) __constant memory__只在定义它的翻译单元中可用，除非您打开单独编译(使用CUDA 5.0或更高版本)。

我仍然不完全明白这对我的情况意味着什么。

假设我不能打开单独编译，有没有办法满足我的需求？我应该把常量内存数组的定义放在哪里？如果我把它放在我的头中，在许多翻译单元中都是#included，会怎么样呢？

假设我可以打开单独的编译，我是否应该在头文件中将我的__constant__内存数组声明为extern，并将定义放在一个源文件中(例如GlobalParameters.cu)？

Answer 1

使常量内存可用于除声明常量的转换单元之外的转换单元的一种方法是调用cudaGetSymbolAddress()并使指针可用于其他函数。

这种技术在某种程度上是在玩火，因为如果在没有适当的屏障和同步的情况下使用指针写入内存，您可能会遇到常量内存和全局内存之间缺乏一致性的问题。

此外，如果使用此方法，您可能无法获得常量内存的全部性能优势。在SM 2.x和更高版本的硬件上，这应该不是那么正确-反汇编目标代码，并确保编译器发出“加载统一”指令。

Answer 2

下面的示例假设可以使用单独的编译。在这种情况下，下面的示例展示了如何使用extern在不同的编译单元中使用常量内存。

文件kernel.cu

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#include "Utilities.cuh"

__constant__ int    N_GPU;
__constant__ float  a_GPU;

__global__ void printKernel();

int main()
{
    const int       N = 5;

    const float     a = 10.466;

    gpuErrchk(cudaMemcpyToSymbol(N_GPU, &N, sizeof(int)));
    gpuErrchk(cudaMemcpyToSymbol(a_GPU, &a, sizeof(float)));

    printKernel << <1, 1 >> > ();
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    return 0;
}

文件otherCompilationUnit.cu

#include <stdio.h>

extern __constant__ int     N_GPU;
extern __constant__ float   a_GPU;

__global__ void printKernel() {

    printf("N = %i; a = %f\n", N_GPU, a_GPU);

}

Answer 3

不，如果不使用单独的编译，就不可能在多个.cu文件中使用相同的常量内存，即只声明一次。

在我看来，有两种方法可以解决这个问题。

第一个是在一个.cu文件中实现所有内核。因此，您将得到的缺点是，此文件将变得非常大，并且概述不好。

第二种方法是在每个.cu文件中再次声明常量内存。然后，像我在回答here中描述的那样，使用包装器将每个.cu文件的值复制到常量内存中。缺点是，您必须确保不会忘记复制单个.cu文件中的值，并且必须检查是否不会在总可用常量内存的限制下运行。