ROCm/Thrust

Query : How to invoke a non device function from a device function

Closed this issue · 2 comments

In the below code we are trying to invoke a non-device function from device function.

/* host function */
int my_rand(void)
{
static thrust::default_random_engine rng;
static thrust::uniform_int_distribution dist(0, 9999);
return dist(rng);
}

int main(void)
{
// generate random data on the host
thrust::generate(h_vec.begin(), h_vec.end(), my_rand); /* generate(device function) internally invokes my_rand(host function)*/
}

Called function is host function and has __ host __ as qualifier
Calling function is device function and has __ device __as qualifier

Facing the below issue while compiling an application in Thrust/examples
./thrust/detail/function.h:52:31: error: call from AMP-restricted function to CPU-restricted function

We tried to compile by adding __ device __ keyword while declaring the m_f function variable but the issue persists

Steps to reproduce the above error:
Step 1: Clone thrust repository from https://github.com/ROCmSoftwarePlatform/Thrust.git

Step 2:In Thrust/thrust/detail/function.h line no:51
-#ifdef __ HIP_PLATFORM_NVCC __
return static_cast(m_f(thrust::raw_reference_cast(x)));
-#endif

Step 3:In Thrust/examples/sum.cpp line:20(Uncomment that line)
+thrust::generate(h_vec.begin(), h_vec.end(), my_rand);

Step 4: Compile using the below command
$ cp sum.cu sum.cpp
$ /opt/rocm/hip/bin/hipcc sum.cpp -I ../ -o sum.cpp.out

full_log.txt

@sriharikarnam considering that this is actively against the CUDA specification as I understand it, I am rather certain that is not what you want to do. In fact, the error that you are getting is EXACTLY expected behaviour, please read Appendix B from here: http://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf, specially B.1.3, which specifically states (paraphrasing) "a host function [...] is callable from the host only". What are you actually trying to achieve? Either the example is bad or you are asking about how to do something which is illegal, the answer to that inquiry being you do not do it.

Thanks Alex we rectified the issue earlier by guarding the host function to be called from host code .
We can close this issue.