How to create and initialize constant a glm::vec3 or float3 array?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I’m using OptiX to render a scene and in one of the shaders I’d need a Poisson disc to sample my light source. So I tried to define the array containing the data in a separate header like so:

__constant__ glm::vec3 samples_8[] =
{
	glm::vec3(0.8308972f, 0.4797409f, 0.f),
	glm::vec3(0.6543578f, -0.48124874f, 0.f),
	glm::vec3(0.11993613f, 0.96366435f, 0.f),
	glm::vec3(0.045692872f, -0.032700025f, 0.f),
	glm::vec3(-0.13769059f, -0.98953694f, 0.f),
	glm::vec3(-0.64136887f, 0.7488851f, 0.f),
	glm::vec3(-0.77864605f, -0.62725335f, 0.f),
	glm::vec3(-0.7274978f, 0.06506414f, 0.f)
};

which failed, telling me that dynamic initialization is not supported for a __constant__ variable. So, next thing I tried was to create a

__constant__ glm::vec3 samples_8[8];

along with the array defined as standard array as it would be used in C/C++ and tried to copy the data using

cudaMemcpyToSymbol(optix::on_device::samples_8, optix::data::samples_8, sizeof(optix::data::samples_8))

which resulted in a cudaErrorInvalidSymbol (invalid device symbol) error. I then read that the __constant__ variables need to be placed in the same *.cu file, which I tried next in the same way as in the first approach (same result, though).

I know this question might appear silly, but it keeps me changing the entire rendering code over and over again for two days now: What is the easiest way to get that array to work in a __global__ OptiX shader? I also know this is sort of mixed up between Cuda and OptiX but I guess this is rather a Cuda issue.

Thank you very much for your help

2

Assuming glm::vec3 is a POD type, the following is worth a try:

__constant__ glm::vec3 samples_8[8] =
{
    {0.8308972f, 0.4797409f, 0.f},
    {0.6543578f, -0.48124874f, 0.f},
    {0.11993613f, 0.96366435f, 0.f},
    {0.045692872f, -0.032700025f, 0.f},
    {-0.13769059f, -0.98953694f, 0.f},
    {-0.64136887f, 0.7488851f, 0.f},
    {-0.77864605f, -0.62725335f, 0.f},
    {-0.7274978f, 0.06506414f, 0.f}
};
3

Hello,

thanks for answering. I haven’t considered that, yet. And although a glm::vec3 can be initialized like that, it, unfortunately, doesn’t solve the “dynamic initialization is not supported for a __constant__ variable” issue.

You’d think it’s easily done but turns out to be a boss fight…

4

Dynamic initialization is not supported for __constant__ data, as the error message says. That is by design. The subject line of this question only refers to initialization.

Below is a worked example of dynamically updating statically initialized __constant__ data. The output of the program should look similar to this:

original statically initialized __constant__ data
 8.30897212e-01  4.79740888e-01  0.00000000e+00
 6.54357791e-01 -4.81248736e-01  0.00000000e+00
 1.19936131e-01  9.63664353e-01  0.00000000e+00
 4.56928723e-02 -3.27000245e-02  0.00000000e+00
-1.37690589e-01 -9.89536941e-01  0.00000000e+00
-6.41368866e-01  7.48885095e-01  0.00000000e+00
-7.78646052e-01 -6.27253354e-01  0.00000000e+00
-7.27497816e-01  6.50641397e-02  0.00000000e+00
dynamically updating second half of __constant__ data
 8.30897212e-01  4.79740888e-01  0.00000000e+00
 6.54357791e-01 -4.81248736e-01  0.00000000e+00
 1.19936131e-01  9.63664353e-01  0.00000000e+00
 4.56928723e-02 -3.27000245e-02  0.00000000e+00
 4.09999990e+00  4.19999981e+00  4.30000019e+00
 5.09999990e+00  5.19999981e+00  5.30000019e+00
 6.09999990e+00  6.19999981e+00  6.30000019e+00
 7.09999990e+00  7.19999981e+00  7.30000019e+00
dynamically updating first half __constant__ data
 1.00000001e-01  2.00000003e-01  3.00000012e-01
 1.10000002e+00  1.20000005e+00  1.29999995e+00
 2.09999990e+00  2.20000005e+00  2.29999995e+00
 3.09999990e+00  3.20000005e+00  3.29999995e+00
 4.09999990e+00  4.19999981e+00  4.30000019e+00
 5.09999990e+00  5.19999981e+00  5.30000019e+00
 6.09999990e+00  6.19999981e+00  6.30000019e+00
 7.09999990e+00  7.19999981e+00  7.30000019e+00

#include <stdio.h>

#define N (8)

typedef struct my_vec3 {
    float x;
    float y;
    float z;
} my_vec3;

__constant__ my_vec3 samples_8[N] =
{
    {0.8308972f, 0.4797409f, 0.f},
    {0.6543578f, -0.48124874f, 0.f},
    {0.11993613f, 0.96366435f, 0.f},
    {0.045692872f, -0.032700025f, 0.f},
    {-0.13769059f, -0.98953694f, 0.f},
    {-0.64136887f, 0.7488851f, 0.f},
    {-0.77864605f, -0.62725335f, 0.f},
    {-0.7274978f, 0.06506414f, 0.f}
};

__global__ void print_samples (void)
{
    for (int i = 0; i < N; i++) {
        printf ("% 15.8e %15.8e %15.8e\n", 
                samples_8[i].x, samples_8[i].y, samples_8[i].z);
    }
} 

void update_samples (my_vec3 *new_samples, size_t size, int offset)
{
    cudaMemcpyToSymbol (samples_8, new_samples, size, offset, cudaMemcpyHostToDevice);
}

int main (void)
{
    my_vec3 new_samples_8[N] = {
        {0.1f, 0.2f, 0.3f},
        {1.1f, 1.2f, 1.3f},
        {2.1f, 2.2f, 2.3f},
        {3.1f, 3.2f, 3.3f},
        {4.1f, 4.2f, 4.3f},
        {5.1f, 5.2f, 5.3f},
        {6.1f, 6.2f, 6.3f},
        {7.1f, 7.2f, 7.3f},

    };    

    printf ("original statically initialized __constant__ data\n");
    print_samples<<<1,1>>>();
    cudaDeviceSynchronize();
    printf ("dynamically updating second half of __constant__ data\n");
    update_samples (new_samples_8 + (N/2), (N/2)*sizeof(my_vec3), (N/2)*sizeof(my_vec3));
    print_samples<<<1,1>>>();
    cudaDeviceSynchronize();
    printf ("dynamically updating first half __constant__ data\n");
    update_samples (new_samples_8, (N/2)*sizeof(my_vec3), 0);
    print_samples<<<1,1>>>();
    return EXIT_SUCCESS;
}

Note: Last I checked the documentation (long time ago), __constant__ has file scope linkage. That means the data and the function that updates it via cuda MemcpyToSymbol() must reside in the same .cu source file. The updating function then can be called from anywhere passing a pointer to the new data.

5

Thank you for your working example.

It seems like glm::vec3 are always dynamically initialized. I tried float3 vectors before, but I initialized them with make_float3, which didn’t work either.

Though, your suggestion to initialize the vector elements with {...} got me thinking… Defining the arrays like so:

__constant__ float3 samples_8[] =
{
	{0.8308972f, 0.4797409f, 0.f},
	{0.6543578f, -0.48124874f, 0.f},
	{0.11993613f, 0.96366435f, 0.f},
	{0.045692872f, -0.032700025f, 0.f},
	{-0.13769059f, -0.98953694f, 0.f},
	{-0.64136887f, 0.7488851f, 0.f},
	{-0.77864605f, -0.62725335f, 0.f},
	{-0.7274978f, 0.06506414f, 0.f}
};

and converting the entries to a glm::vec3 when needed was the solution.

Hell… three days for such an easy thing ;_;

6

my guess is that glm::vec3 has a constructor. That’s going to trip the dynamic initialization error.

The CUDA provided vector types don’t have constructors.

7

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.