diff --git a/sycl/test-e2e/DeviceLib/built-ins/ext_ldg.cpp b/sycl/test-e2e/DeviceLib/built-ins/ext_ldg.cpp new file mode 100644 index 0000000000000..29f51281a0490 --- /dev/null +++ b/sycl/test-e2e/DeviceLib/built-ins/ext_ldg.cpp @@ -0,0 +1,111 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %t.out + +// Checks that the sycl::ext::oneapi::experimental::cuda::ldg builtins are +// returning the correct values. + +#include +#include + +using namespace sycl::ext::oneapi::experimental::cuda; +using namespace sycl::ext::oneapi::experimental; +using namespace sycl; + +template class KernelName; + +template bool checkEqual(vec A, vec B) { + return A.x() == B.x() && A.y() == B.y(); +} + +template bool checkEqual(vec A, vec B) { + return A.x() == B.x() && A.y() == B.y() && A.z() == B.z() && A.w() == B.w(); +} + +template void test(sycl::queue &q) { + + T a_loc; + T b_loc; + + a_loc = 2; + b_loc = 3; + + T *A = malloc_device(1, q); + T *B = malloc_device(1, q); + T *C = malloc_device(1, q); + + q.memcpy(A, &a_loc, sizeof(T)); + q.memcpy(B, &b_loc, sizeof(T)); + q.wait(); + + q.submit([=](sycl::handler &h) { + h.parallel_for>(range<1>(1), [=](auto i) { + auto cacheA = ldg(&A[0]); + auto cacheB = ldg(&B[0]); + C[0] = cacheA + cacheB; + }); + }); + + T dev_result; + q.wait(); + + q.memcpy(&dev_result, C, sizeof(T)).wait(); + if constexpr (std::is_same_v || std::is_same_v || + std::is_same_v || std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v || std::is_same_v) { + + assert(dev_result == a_loc + b_loc); + } else { + assert(checkEqual(dev_result, a_loc + b_loc)); + } + + free(A, q); + free(B, q); + free(C, q); +} + +int main() { + queue q; + + test(q); + test(q); + test(q); + test(q); + test(q); + + test(q); + test(q); + test(q); + + test(q); + + test(q); + test(q); + test(q); + + test(q); + test(q); + test(q); + + test(q); + test(q); + test(q); + + test(q); + test(q); + test(q); + + test(q); + test(q); + test(q); + + if (q.get_device().has(sycl::aspect::fp64)) { + test(q); + test(q); + } + return 0; +}