In my program I want to use a structure containing constant variables and keep it on device all long as the program executes to completion.
I have several header files containing the declaration of 'global' functions and their respective '.cu' files for their definitions. I kept this scheme because it helps me contain similar code in one place. e.g. all the 'device' functions required to complete 'KERNEL_1' are separated from those 'device' functions required to complete 'KERNEL_2' along with kernels definitions.
I had no problems with this scheme during compilation and linking. Until I encountered constant variables. I want to use the same constant variable through all kernels and device functions but it doesn't seem to work.
##########################################################################
CODE EXAMPLE
###########################################################################
filename: 'common.h'
--------------------------------------------------------------------------
typedef struct {
double height;
double weight;
int age;
} __CONSTANTS;
__constant__ __CONSTANTS d_const;
---------------------------------------------------------------------------
filename: main.cu
---------------------------------------------------------------------------
#include "common.h"
#include "gpukernels.h"
int main(int argc, char **argv) {
__CONSTANTS T;
T.height = 1.79;
T.weight = 73.2;
T.age = 26;
cudaMemcpyToSymbol(d_consts, &T, sizeof(__CONSTANTS));
test_kernel <<< 1, 16 >>>();
cudaDeviceSynchronize();
}
---------------------------------------------------------------------------
filename: gpukernels.h
---------------------------------------------------------------------------
__global__ void test_kernel();
---------------------------------------------------------------------------
filename: gpukernels.cu
---------------------------------------------------------------------------
#include <stdio.h>
#include "gpukernels.h"
#include "common.h"
__global__ void test_kernel() {
printf("Id: %d, height: %f, weight: %f\n", threadIdx.x, d_const.height, d_const.weight);
}
When I execute this code, the kernel executes, displays the thread ids, but the constant values are displayed as zeros. How can I fix this?
