Defining global variables on the host and device at once?

Is there a way I can define a constant as such

int c = 500;

and have it be defined as global for both the host and device at once? Or is the only way to define on the host first and then use a memory copy function?

Yes, just pass the constant as a parameter to the kernel.

An idiom I’ve been using lately is aggregating all constants – values, constant pointers, etc. – into a single constant structure and passing that constant “environment” down through the kernel’s device functions as needed (while being careful to maintain its const’ness).

The compiler does a great job recognizing that these values and pointers reside in the constant memory space. Works well for sm_20+.

I’m a little bit uncertain of that idiom because where do I define what the constants are in value? What I’ve tried so far, for one small structure is the following

Out in host global space

struct matl1
{
	static const double cond;
};

const double matl1::cond = 420.5;

Then inside of main()

matl1 * h_matl1 = (matl1*)malloc(sizeof(matl1));
matl1 * d_matl1;
cudaMalloc((void**)&d_matl1, sizeof(matl1));
cudaMemcpy(d_matl1, h_matl1, sizeof(matl1), cudaMemcpyHostToDevice);
kernel<<<1,1>>>(d_matl1,...);

Then inside of kernel()

__global__ void mainCalc(matl1* d_matl1,...)
{
double cond = d_matl1->cond;
}

And I get the following error:

error : identifier “matl1::cond” is undefined in device code

As a quick test, if I do the following on the host in main()

cout << h_matl1->cond << endl;

It shows me the correct output of 420.5. I am not sure why this isn’t making it into the device.

Anyone with any ideas, or fixes? I haven’t gotten anywhere in the past day. :(

Actually I think this is the same problem allanmac helped me solved a few weeks ago. You might want to check out

https://devtalk.nvidia.com/default/topic/628417/help-with-pointers-inside-structs/#3991261

Alternatively I have copied his reply below

“I think the issue is that you’re passing “hostStruct” by value. Passing the struct by reference should fix it:”

__host__ void fcreate( testStruct& hostStruct ) {
...
}

I don’t think that was my. I can’t even define it like that. That is a host function and I am working with the device.

I’m going to try going the other way. Start by defining everything on the device, and then hope it works copying it over to the host, as opposed to my initial failure of host to device.

Just an update, from device to host works. So I guess I’ll be working with that. That is, defining all of my constants in the device code, and passing them back to the host, as opposed to vice versa.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <stdio.h>

using namespace std;

typedef struct {
    double cond, rho, cp;
} matl1;

__device__  void foo(matl1 *d_matl1) {
    d_matl1->cond = 420.5;
    d_matl1->rho = 8960.0;
	d_matl1->cp = 390.0;
}

__global__ void myKernel(matl1 *d_matl1){
	foo(d_matl1);
	printf("Device is %.2f\n",d_matl1->cond);
}

int main() {

	matl1 *h_matl1, *d_matl1;

	h_matl1 = (matl1*)malloc(sizeof(matl1));
    cudaMalloc((void**)&d_matl1, sizeof(matl1));
	
	myKernel  <<<1,1 >>>(d_matl1);
    cudaThreadSynchronize();

	cudaMemcpy(h_matl1, d_matl1, sizeof(matl1), cudaMemcpyDeviceToHost);

	printf("Host is %.2f\n",h_matl1->cond);
	cin.ignore();
    return 0;
}

I’m sure you could have replaced the host with a global in my example.

const int c = 500;

defined at global scope, should be visible to both host and device code defined in that compilation unit.