将静态声明的 __constant__ 变量作为内核参数（CUDA）传递答案

【问题标题】：Pass statically declared __constant__ variable as kernel parameter (CUDA)将静态声明的 __constant__ 变量作为内核参数（CUDA）传递
【发布时间】：2020-04-01 15:15:35
【问题描述】：

我正在尝试将静态分配的 __constant__ 变量作为内核参数传递，但似乎有问题。内核表现为变量未初始化。但是，我仍然可以将变量作为全局变量访问（来自同一个内核）。从主机修改变量值有效，但同样，只有当我从全局范围访问变量时。这是一个例子：

#include<cuda_runtime_api.h>
#include<cuda_runtime.h>
#include<stdio.h>

__constant__ float constant_a=1.12345;
__constant__ float constant_b;

__global__ void test_values(float a, float b) {
    printf("Device code: constant_a = %f, constant_b = %f\n", a, b);
    printf("Device code: constant_a = %f, constant_b = %f\n\n", constant_a, constant_b);
}

int main() {
    test_values<<<1, 1>>>(constant_a, constant_b);
    cudaDeviceSynchronize();

    const float h_const_a = 1;
    const float h_const_b = 2;

    cudaMemcpyToSymbol(constant_a, &h_const_a, sizeof(float));
    test_values<<<1, 1>>>(constant_a, constant_b);
    cudaDeviceSynchronize();

    cudaMemcpyToSymbol(constant_b, &h_const_b, sizeof(float));
    test_values<<<1, 1>>>(constant_b, constant_b);
    cudaDeviceSynchronize();
}

内核打印出这个：

Device code: constant_a = 0.000000, constant_b = 0.000000
Device code: constant_a = 1.123450, constant_b = 0.000000

Device code: constant_a = 0.000000, constant_b = 0.000000
Device code: constant_a = 1.000000, constant_b = 0.000000

Device code: constant_a = 0.000000, constant_b = 0.000000
Device code: constant_a = 1.000000, constant_b = 2.000000

如果有人可以解释这一点，如果您能提供信息来源，我也将非常感激。因为我没有在 Nvidia Guide 或其他几本书中找到这些信息。

【问题讨论】：

__constant__ 和 __device__ 变量存在于全局范围内。没有必要将它们作为参数传递，实际上你不应该那样使用它们。只需在您的代码中使用它们，就好像它们在模块中全局可用一样。因为他们是。记录这一点的地方之一是here“可从网格内的所有线程访问”意味着无需您将其作为参数传递即可访问它。
网格中的所有线程都可以访问所有内核参数，因此如果您仍然需要将这些变量作为参数传递，则无需用该措辞指出这一点。那不是本意。目的是当您使用__constant__ 进行装饰时，网格中的所有线程都可以自动访问它们。
@RobertCrovella 非常感谢您的回复！我阅读了文档，但没有发现将它们作为参数传递是错误的，所以我想也许我可以同时使用这两种方法——传递 constant 变量的方式与传递动态分配的设备数组相同。无论如何，我错了，但我仍然认为 Nvidia 应该更明确地将这些信息放在某个地方。谢谢！

标签： cuda

【解决方案1】：

我正在尝试将静态分配的 __constant__ 变量传递为内核参数，但似乎有问题

问题是你不能这样做。不支持。

内核表现为变量未初始化

因为它们没有被初始化。 __constant__ 变量不能在主机代码中访问，除非通过符号操作 API。当您使用这些 API 时，您只是在修改设备内存中的值。没有其他的。与静态声明的设备符号关联的主机端变量仅作为绑定 API 调用的标签存在。设备内存的变化不会反映在主机内存中。唯一的例外是 __managed__ 符号（在支持的平台上）。

从主机修改变量值有效，但同样，只有当我从全局范围访问变量

这是主机代码中 __constant__ 变量唯一有效且受支持的用例。

因为我没有在 Nvidia Guide 或其他几本书中找到此信息。

您没有找到有关如何执行此操作的信息，因为它不存在，那是因为它不受支持。您可以在here 找到支持的内容的简明描述。

【讨论】：