host float constant usage in a kernel in CUDA -

host float constant usage in a kernel in CUDA -

i using cuda 5.0. noticed compiler allow me utilize host-declared int constants within kernels. however, refuses compile kernels utilize host-declared float constants. know reason seeming discrepancy?

for example, next code runs fine is, not compile if final line in kernel uncommented.

#include <cstdio> #include <cuda_runtime.h> static int __constant__ dev_int_constant = 1; static float __constant__ dev_float_constant = 2.0f; static int const hst_int_constant = 3; static float const hst_float_constant = 4.0f; __global__ void uselesskernel(float * val) { *val = 0.0f; // utilize device int , float constants *val += dev_int_constant; *val += dev_float_constant; // utilize host int , float constants *val += hst_int_constant; //*val += hst_float_constant; // won't compile if uncommented } int main(void) { float * d_val; cudamalloc((void **)&d_val, sizeof(float)); uselesskernel<<<1, 1>>>(d_val); cudafree(d_val); }

thank you,

aaron

adding const number in device code ok, adding number stored on host memory in device code not.

every reference of static const int in code can replaced value 3 compiler/optimizer when addr of variable never referenced. in case, #define hst_int_constant 3, , no host memory allocated variable.

but float var, host memory allocated of static const float. since kernel can not access host memory directly, code static const float won't compiled.

for c/c++, int can optimized more aggressively float.

you code runs when comment on can seen bug of cuda c think. static const int host side thing, , should not accessible device directly.

cuda kernel constants host