Using constant memory with separate compilation

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

is there a way to combine constant memory and separate compilation. Seemingly, according to my interpretation of the programming guide, you can not pass pointers obtained for constant memory on/ by the device, to the host, nor pointers for constant memory obtained for constant memory obtained by the host, to the device (programming guide: E.2.3. Pointers) - i.e. direct memory copies of pointers.

Would you be able to use a function that obtains the constant memory addresses via cudaGetSymbolAddress and pass it to other functions and program files via a structure? The function is made to reside in the source file where the constants are used in device kernels, obtains the constant addresses via cudaGetSymbolAddress and shares this with application functions residing in other source files via a structure.

2 
__constant__ int c_markov_write_to_arr;
__constant__ int* c_markov_rfr_coeff_mem_nr;
__constant__ int* c_markov_rfr_coeff_mem_map;
__constant__ double* c_markov_rfr_coeff_mem_prob;

void rfr_opt_set_markov_otf_const(int* flag, RFR* rfr, cudaStream_t* s0)
{
	cudaMemcpyToSymbolAsync(c_markov_write_to_arr,
		flag,
		sizeof(int), 0,
		cudaMemcpyHostToDevice, s0[0]);
}


void rfr_opt_set_markov_mem_const(int** mem_nr, int** mem_map, double** mem_prob,
	RFR* rfr, cudaStream_t* s0)
{
	cudaMemcpyToSymbolAsync(c_markov_rfr_coeff_mem_nr,
		mem_nr,
		sizeof(int*), 0,
		cudaMemcpyHostToDevice, s0[0]);

	cudaMemcpyToSymbolAsync(c_markov_rfr_coeff_mem_map,
		mem_map,
		sizeof(int*), 0,
		cudaMemcpyHostToDevice, s0[0]);

	cudaMemcpyToSymbolAsync(c_markov_rfr_coeff_mem_prob,
		mem_prob,
		sizeof(double*), 0,
		cudaMemcpyHostToDevice, s0[0]);
}

seems to work. the other source files call the functions <rfr_opt_set_markov_otf_const> and <rfr_opt_set_markov_mem_const>

3

A pointer to constant memory should be valid however it is used within a given device context. So it should be possible to pass a pointer around to reference it. You seem to be already indicating this is working for you.

It should not be necessary to pass a pointer around, however. If you properly use the extern keyword, the constant symbol in a given compilation unit can be made to “appear” in another compilation unit (in a device-linked setting) and can be referenced directly (by name) that way.

[url]CUDA constant memory symbols - Stack Overflow