0

I have init function er_const.cu

__device__ __constant__ __extended_range_struct ER_ONE_CONST;
__device__ __constant__ __extended_range_struct ER_HALF_CONST;
__device__ __constant__ __extended_range_struct ER_SERIES_PREC_CONST;

namespace cuda {
    void er_const_init() {
        er_ptr one_host;
        one_host = new __extended_range_struct;
        *one_host = (__extended_range_struct) {1.0, 0};
        gpuErrchk(cudaMemcpyToSymbol(ER_ONE_CONST, one_host, sizeof(__extended_range_struct)));
        gpuErrchk(cudaDeviceSynchronize());

        er_ptr half_host;
        half_host = new __extended_range_struct;
        er_set_d(half_host, 0.5);
        gpuErrchk(cudaMemcpyToSymbol(ER_HALF_CONST, half_host, sizeof(__extended_range_struct)));
        gpuErrchk(cudaDeviceSynchronize());
        er_ptr series_prec_host;
        series_prec_host = new __extended_range_struct;
        er_set_d(series_prec_host, ER_SERIES_PREC);
        gpuErrchk(cudaMemcpyToSymbol(ER_SERIES_PREC_CONST, series_prec_host, sizeof(__extended_range_struct)));
        gpuErrchk(cudaDeviceSynchronize());
        delete one_host;
        delete half_host;
        delete series_prec_host;
    }
}

__attribute__ ((constructor)) void premain()
{
        //cuda::er_const_init();

}

and header file for it er_const.cuh

extern __device__ __constant__ __extended_range_struct ER_ONE_CONST;
extern __device__ __constant__ __extended_range_struct ER_HALF_CONST;
extern __device__ __constant__ __extended_range_struct ER_SERIES_PREC_CONST;

namespace cuda {
    void er_const_init();
}

if i call this function in main() it's work, but if i try create function in er_const.cu wich called before main() i have error: "invalid device symbol"

talonmies
  • 67,081
  • 33
  • 170
  • 244
Alex Kuvaev
  • 37
  • 5
  • 3
    The really short answer is that you can't do this. The long answer is explained in the marked duplicate. There are "invisible" static objects which the CUDA toolchain emits and which must be instantiated before any CUDA code which accesses kernels, global device symbols or textures can successfully run. When you put such code at the same scope as those CUDA objects, there are no guarantees that the CUDA setup code will run before yours, and it can and does fail – talonmies Aug 28 '16 at 14:37
  • CUDA dont have any events in this "invisible" static objects? or any constructors, desctructors? – Alex Kuvaev Aug 29 '16 at 09:46
  • There is a constructor and the registration of several callbacks to boilerplate functions emitted by the compiler which trigger context initialisation and the loading of the GPU binary payload to the device. If you want to see how it works, compile your code with the `--keep` option and have a look at the intermediate code produced by the toolchain for yourself – talonmies Aug 29 '16 at 11:18

0 Answers0