__constant__ で宣言した構造体の変数を初期化する場合は初期化子リストか constexpr で宣言した変数を使用する【CUDA】

CUDA

試した環境

本題

先に試したコードを載せます。

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

#include <stdio.h>

struct Point {
    int x;
    int y;
};

__constant__ Point dev_P0 = { 100, 15 };

Point host_P1_1 = { 10, 25 };
const Point host_P1_2 = { 20, 45 };
constexpr Point host_P1_3 = { 30, 65 };

__constant__ Point dev_P1 = host_P1_3;


__global__ void addConstant(Point* dst, const Point* src) {
    int i = threadIdx.x;
    dst[i].x = src[i].x + dev_P0.x + dev_P1.x;
    dst[i].y = src[i].y + dev_P0.y + dev_P1.y;
}

int main()
{
    const int array_size = 5;
    const Point src[array_size] = { {1, -1}, {2, -2}, {3, -3}, {4, -4}, {5, -5} };
    Point dst[array_size] = { 0 };

    Point* dev_src = nullptr;
    Point* dev_dst = nullptr;

    cudaMalloc((void**)&dev_src, array_size * sizeof(Point));
    cudaMalloc((void**)&dev_dst, array_size * sizeof(Point));

    cudaMemcpy(dev_src, src, array_size * sizeof(Point), cudaMemcpyHostToDevice);

    addConstant <<<1, array_size >>> (dev_dst, dev_src);

    cudaMemcpy(dst, dev_dst, array_size * sizeof(Point), cudaMemcpyDeviceToHost);

    printf("{ {1,-1}, {2,-2}, {3,-3}, {4,-4}, {5,-5} } + P0 + P1 = { {%d,%d}, {%d,%d}, {%d,%d}, {%d,%d}, {%d,%d} }\n",
        dst[0].x, dst[0].y, dst[1].x, dst[1].y, dst[2].x, dst[2].y, dst[3].x, dst[3].y, dst[4].x, dst[4].y);

    cudaFree(dev_dst);
    cudaFree(dev_src);

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaError_t cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

__constant__で宣言した構造体（上記コードではstruct Point）の変数を初期化する場合は、上記コードのdev_P0のように初期化子リストを使うかdev_P1のようにconstexprで宣言した変数を使用します。 __constant__ Point dev_P1 = host_P1_3;はビルドが通りますが、__constant__ Point dev_P1 = host_P1_1;や__constant__ Point dev_P1 = host_P1_2;にするとビルドエラーになります。
エラーメッセージは以下の通りです。

dynamic initialization is not supported for a __constant__ variable