Linking cuda code with SYCL ? #3627
dosarudaniel
started this conversation in
General
Replies: 1 comment 9 replies
-
|
I don't have an answer to this question and I'm not sure that we have such functionality, so I will ping a few people here first to see if they know something: @AGindinson, @steffenlarsen. Additionally, a few directions where to look: You can try to do the operation you want manually: to do so, you can launch your |
Beta Was this translation helpful? Give feedback.
9 replies
-
|
I was trying to generate a bytecode for CUDA part and SYCL device code and then to link it, but it doesn't work since different target triples are used, e.g. However, the output file is generated. I assume it will not work or it is possible to compile it further and use with the host code compilation? @AGindinson, @steffenlarsen, @alexbatashev any ideas? |
Beta Was this translation helpful? Give feedback.
-
|
I wondered if you were ever able to solve this? I have a similar example which would need this kind of solution (to use calls to VecGeom code, which only supports CUDA as far as I know, in an application that would use SYCL). Cheers, Mark |
Beta Was this translation helpful? Give feedback.
-
|
Hi @shefmarkh, I did not solve it, I think it is one of the limitations of the open source compiler.. Here is the full story: DevSummit presentation If you will find a workaround let me know as I am still interested in it. The external library has CUDA entangled in types and interfaces and it is hard to migrate it to SYCL.. |
Beta Was this translation helpful? Give feedback.
-
|
Sorry. I completely missed this discussion. So I think @ivorobts hit the nail on the head. Currently I think the best way to achieve what you're after is to add the bitcode generated using clang CUDA to the device-code link step. The following works on my machine. I've changed the code slightly. The use of field.cu: __host__ __device__ double stepInField(double kinE, double mass, int charge)
{
return(kinE*mass*charge);
}test12.cpp: #include <CL/sycl.hpp>
extern SYCL_EXTERNAL double stepInField(double kinE, double mass, int charge); // Check field.cu file
void kernel(double *step)
{
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
*step = stepInField(6.0f, 7.0f,1);
#endif
}
int main(void)
{
sycl::default_selector device_selector;
sycl::queue q_ct1(device_selector);
std::cout << "Running on "
<< q_ct1.get_device().get_info<cl::sycl::info::device::name>()
<< "\n";
double *d_dev_ptr;
d_dev_ptr = sycl::malloc_device<double>(1, q_ct1);
q_ct1.submit([&](sycl::handler &cgh) {
cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, 1),
sycl::range<3>(1, 1, 1)),
[=](sycl::nd_item<3> item_ct1) {
kernel(d_dev_ptr);
});
}).wait();
double d_dev;
q_ct1.memcpy(&d_dev, d_dev_ptr, sizeof(double)).wait();
std::cout << " device: " << d_dev << std::endl;
sycl::free(d_dev_ptr, q_ct1);
}To build these I am using DPC++ HEAD e8f2166. First step, as @ivorobts suggested, is to compile the CUDA code to LLVM bitcode: This gives us the To find the device-code linking step we use Of interest here is the step using This generates the I hope that helps! Sorry again for not responding earlier, @dosarudaniel . I tried to see if there was a better way to specify additional device-code link files in the driver but I couldn't find anything. Maybe @mdtoguchi and @AGindinson knows if there is. |
Beta Was this translation helpful? Give feedback.
-
|
ok thanks! Will take a look at that git hash then :) Thanks also for the slides @dosarudaniel , I found them very useful. Cheers, Mark |
Beta Was this translation helpful? Give feedback.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Uh oh!
There was an error while loading. Please reload this page.
Uh oh!
There was an error while loading. Please reload this page.
-
Hi,
I am working on a simple example called test12.cpp to mimic on a smaller scale the project I am working on (converting the AdePT project example9 written in cuda to dpcpp).
Currently, in test12.cpp, I have the stepInField function defined in the same file and used in the kernel, but would it be possible to place it in another file (like field.cu) and compile this field.cu file with nvcc and link it with the main SYCL program (test12) where the stepInField is called in the kernel ?
Basically, for the project I am working on I basically need to call a CUDA function (like fieldInStep)) from a SYCL kernel (like this one)
Can this be done with clang++ ?
Thanks,
Daniel
Beta Was this translation helpful? Give feedback.
All reactions