// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --functions "ldg" --include-generated-funcs --version 4
// REQUIRES: cuda

// RUN: %clangxx -fsycl-device-only -fsycl-targets=nvptx64-nvidia-cuda -Xclang -fnative-half-type -S -Xclang -emit-llvm %s -o -| FileCheck %s --check-prefixes=CHECK-OPAQUE

#include <sycl/ext/oneapi/experimental/cuda/builtins.hpp>
#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::oneapi::experimental::cuda;
using namespace sycl::ext::oneapi::experimental;

// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL half
sycl::ext::oneapi::experimental::cuda::ldg(const half *);
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL float
sycl::ext::oneapi::experimental::cuda::ldg(const float *);
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL double
sycl::ext::oneapi::experimental::cuda::ldg(const double *);

// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<half, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<half, 2> *);
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<half, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<half, 3> *);
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load half, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<half, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<half, 4> *);
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<float, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<float, 2> *);
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<float, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<float, 3> *);
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load float, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<float, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<float, 4> *);
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<double, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<double, 2> *);
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<double, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<double, 3> *);
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load double, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<double, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<double, 4> *);

// Unsigned variants are identical to signed variants, but this leads to
// correct behavior.

// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL char
sycl::ext::oneapi::experimental::cuda::ldg(const char *);
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL signed char
sycl::ext::oneapi::experimental::cuda::ldg(const signed char *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL short
sycl::ext::oneapi::experimental::cuda::ldg(const short *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL int
sycl::ext::oneapi::experimental::cuda::ldg(const int *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL long
sycl::ext::oneapi::experimental::cuda::ldg(const long *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL long long
sycl::ext::oneapi::experimental::cuda::ldg(const long long *);

// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL unsigned char
sycl::ext::oneapi::experimental::cuda::ldg(const unsigned char *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL unsigned short
sycl::ext::oneapi::experimental::cuda::ldg(const unsigned short *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL unsigned int
sycl::ext::oneapi::experimental::cuda::ldg(const unsigned int *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL unsigned long
sycl::ext::oneapi::experimental::cuda::ldg(const unsigned long *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
template SYCL_EXTERNAL unsigned long long
sycl::ext::oneapi::experimental::cuda::ldg(const unsigned long long *);

// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<char, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<char, 2> *);
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<char, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<char, 3> *);
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<signed char, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<signed char, 2> *);
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<signed char, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<signed char, 3> *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<short, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<short, 2> *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<short, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<short, 3> *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<int, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<int, 2> *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<int, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<int, 3> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<long, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<long, 2> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<long, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<long, 3> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<long long, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<long long, 2> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<long long, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<long long, 3> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<long, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<long, 4> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<long long, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<long long, 4> *);

// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned char, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned char, 2> *);
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned char, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned char, 3> *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned short, 2>
sycl::ext::oneapi::experimental::cuda::ldg(
    const sycl::vec<unsigned short, 2> *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned short, 3>
sycl::ext::oneapi::experimental::cuda::ldg(
    const sycl::vec<unsigned short, 3> *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned int, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned int, 2> *);

// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned int, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned int, 3> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned long, 2>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned long, 2> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned long, 3>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned long, 3> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned long long, 2>
sycl::ext::oneapi::experimental::cuda::ldg(
    const sycl::vec<unsigned long long, 2> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}!invariant.load
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned long long, 3>
sycl::ext::oneapi::experimental::cuda::ldg(
    const sycl::vec<unsigned long long, 3> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned long, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned long, 4> *);
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i64, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned long long, 4>
sycl::ext::oneapi::experimental::cuda::ldg(
    const sycl::vec<unsigned long long, 4> *);

// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<char, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<char, 4> *);
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<signed char, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<signed char, 4> *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<short, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<short, 4> *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<int, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<int, 4> *);

// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i8, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned char, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned char, 4> *);
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i16, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned short, 4>
sycl::ext::oneapi::experimental::cuda::ldg(
    const sycl::vec<unsigned short, 4> *);
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
// CHECK-OPAQUE: load i32, ptr addrspace(1) %{{.*}}
template SYCL_EXTERNAL sycl::vec<unsigned int, 4>
sycl::ext::oneapi::experimental::cuda::ldg(const sycl::vec<unsigned int, 4> *);