Skip to content

Instantly share code, notes, and snippets.

@al42and
Created February 25, 2021 11:24
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save al42and/f85e5e11a3cd6cf448a4e718f2489f73 to your computer and use it in GitHub Desktop.
Save al42and/f85e5e11a3cd6cf448a4e718f2489f73 to your computer and use it in GitHub Desktop.
#include <CL/sycl.hpp>
#include <iostream>
#ifdef USE_NATIVE
static inline float mySqrt(float x) {
#ifdef SYCL_DEVICE_ONLY
return cl::sycl::native::sqrt(x);
#else
return cl::sycl::sqrt(x);
#endif
}
#else
static inline float mySqrt(float x) { return cl::sycl::sqrt(x); }
#endif
int main() {
cl::sycl::queue q{cl::sycl::gpu_selector{}};
const cl::sycl::nd_range<1> range{32, 32};
cl::sycl::buffer<float, 1> buf(32);
q.submit([&](cl::sycl::handler &cgh) {
auto acc = buf.get_access<cl::sycl::access::mode::discard_write>(cgh);
cgh.parallel_for<class KernelName>(
range, [=](cl::sycl::nd_item<1> itemIdx) {
float x = itemIdx.get_global_linear_id();
x = mySqrt(x);
acc[itemIdx.get_global_id()] = x;
});
}).wait_and_throw();
std::cout << "Testing done" << std::endl;
return 0;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment