Subversion Repositories QNX 8.QNX8 LLVM/Clang compiler suite

/*===--- __clang_cuda_texture_intrinsics.h - Device-side texture support ---===

 *

 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.

 * See https://llvm.org/LICENSE.txt for license information.

 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

 *

 *===-----------------------------------------------------------------------===

 *

 * This header provides in-header implmentations for NVCC's built-in

 * __nv_tex_surf_handler() which is used by CUDA's texture-related headers.  The

 * built-in is unusual as it's actually a set of function overloads that use the

 * first string literal argument as one of the overload parameters.

 */

#ifndef __CLANG_CUDA_TEXTURE_INTRINSICS_H__

#define __CLANG_CUDA_TEXTURE_INTRINSICS_H__

#ifndef __CUDA__

#error "This file is for CUDA compilation only."

#endif

// __nv_tex_surf_handler() provided by this header as a macro.

#define __nv_tex_surf_handler(__op, __ptr, ...)                                \

  ::__cuda_tex::__tex_fetch<                                                   \

      ::__cuda_tex::__Tag<::__cuda_tex::__tex_op_hash(__op)>>(__ptr,           \

                                                              __VA_ARGS__)

#pragma push_macro("__ASM_OUT")

#pragma push_macro("__ASM_OUTP")

#pragma push_macro("__Args")

#pragma push_macro("__ID")

#pragma push_macro("__IDV")

#pragma push_macro("__IMPL_2DGATHER")

#pragma push_macro("__IMPL_ALIAS")

#pragma push_macro("__IMPL_ALIASI")

#pragma push_macro("__IMPL_F1")

#pragma push_macro("__IMPL_F3")

#pragma push_macro("__IMPL_F3N")

#pragma push_macro("__IMPL_F3S")

#pragma push_macro("__IMPL_S")

#pragma push_macro("__IMPL_S3")

#pragma push_macro("__IMPL_S3I")

#pragma push_macro("__IMPL_S3N")

#pragma push_macro("__IMPL_S3NI")

#pragma push_macro("__IMPL_S3S")

#pragma push_macro("__IMPL_S3SI")

#pragma push_macro("__IMPL_SI")

#pragma push_macro("__L")

#pragma push_macro("__STRIP_PARENS")

// Put all functions into anonymous namespace so they have internal linkage.

// The device-only function here must be internal in order to avoid ODR

// violations in case they are used from the files compiled with

// -fgpu-rdc. E.g. a library and an app using it may be built with a different

// version of this header file.

namespace {

// Put the implmentation into its own namespace so we don't pollute the TU.

namespace __cuda_tex {

// First, we need a perfect hash function and a few constexpr helper functions

// for converting a string literal into a numeric value which can be used to

// parametrize a template. We can not use string literals for that as that would

// require C++20.

//

// The hash function was generated with 'gperf' and then manually converted into

// its constexpr equivalent.

//

// NOTE: the perfect hashing scheme comes with inherent self-test. If the hash

// function has a collision for any of the texture operations, the compilation

// will fail due to an attempt to redefine a tag with the same value. If the

// header compiles, then the hash function is good enough for the job.

constexpr int __tex_len(const char *s) {

  return (s[0] == 0)    ? 0

         : (s[1] == 0)  ? 1

         : (s[2] == 0)  ? 2

         : (s[3] == 0)  ? 3

         : (s[4] == 0)  ? 4

         : (s[5] == 0)  ? 5

         : (s[6] == 0)  ? 6

         : (s[7] == 0)  ? 7

         : (s[8] == 0)  ? 8

         : (s[9] == 0)  ? 9

         : (s[10] == 0) ? 10

         : (s[11] == 0) ? 11

         : (s[12] == 0) ? 12

         : (s[13] == 0) ? 13

         : (s[14] == 0) ? 14

         : (s[15] == 0) ? 15

         : (s[16] == 0) ? 16

         : (s[17] == 0) ? 17

         : (s[18] == 0) ? 18

         : (s[19] == 0) ? 19

         : (s[20] == 0) ? 20

         : (s[21] == 0) ? 21

         : (s[22] == 0) ? 22

         : (s[23] == 0) ? 23

         : (s[24] == 0) ? 24

         : (s[25] == 0) ? 25

         : (s[26] == 0) ? 26

         : (s[27] == 0) ? 27

         : (s[28] == 0) ? 28

         : (s[29] == 0) ? 29

         : (s[30] == 0) ? 30

         : (s[31] == 0) ? 31

                        : 32;

}

constexpr int __tex_hash_map(int c) {

  return (c == 49)    ? 10

         : (c == 50)  ? 0

         : (c == 51)  ? 100

         : (c == 52)  ? 30

         : (c == 67)  ? 10

         : (c == 68)  ? 0

         : (c == 69)  ? 25

         : (c == 72)  ? 70

         : (c == 77)  ? 0

         : (c == 96)  ? 44

         : (c == 99)  ? 10

         : (c == 100) ? 5

         : (c == 101) ? 60

         : (c == 102) ? 40

         : (c == 103) ? 70

         : (c == 104) ? 25

         : (c == 112) ? 0

         : (c == 114) ? 45

         : (c == 117) ? 5

         : (c == 118) ? 85

         : (c == 120) ? 20

                      : 225;

}

constexpr int __tex_op_hash(const char *str) {

  return __tex_len(str) + __tex_hash_map(str[7] + 1) + __tex_hash_map(str[6]) +

         __tex_hash_map(str[5]) + __tex_hash_map(str[__tex_len(str) - 1]);

}

// Tag type to identify particular texture operation.

template <int N> struct __Tag;

#define __ID(__op) __Tag<__tex_op_hash(__op)>

// Tags for variants of particular operation. E.g. tex2Dgather can translate

// into 4 different instructions.

#define __IDV(__op, __variant)                                                 \

  __Tag<10000 + __tex_op_hash(__op) * 100 + __variant>

// Helper classes for figuring out key data types for derived types.

// E.g. char2 has __base_t = char, __fetch_t = char4

template <class> struct __TypeInfoT;

// Type info for the fundamental types.

template <> struct __TypeInfoT<float> {

  using __base_t = float;

  using __fetch_t = float4;

};

template <> struct __TypeInfoT<char> {

  using __base_t = char;

  using __fetch_t = int4;

};

template <> struct __TypeInfoT<signed char> {

  using __base_t = signed char;

  using __fetch_t = int4;

};

template <> struct __TypeInfoT<unsigned char> {

  using __base_t = unsigned char;

  using __fetch_t = uint4;

};

template <> struct __TypeInfoT<short> {

  using __base_t = short;

  using __fetch_t = int4;

};

template <> struct __TypeInfoT<unsigned short> {

  using __base_t = unsigned short;

  using __fetch_t = uint4;

};

template <> struct __TypeInfoT<int> {

  using __base_t = int;

  using __fetch_t = int4;

};

template <> struct __TypeInfoT<unsigned int> {

  using __base_t = unsigned int;

  using __fetch_t = uint4;

};

// Derived base/fetch types for N-element vectors.

template <class __T> struct __TypeInfoT {

  using __base_t = decltype(__T::x);

  using __fetch_t = typename __TypeInfoT<__base_t>::__fetch_t;

};

// Classes that implement specific texture ops.

template <class __op> struct __tex_fetch_v4;

// Helper macros to strip parens from a macro argument.

#define __Args(...) __VA_ARGS__

#define __STRIP_PARENS(__X) __X

#define __L(__X) __STRIP_PARENS(__Args __X)

// Construct inline assembly output args.

// Results are stored in a temp var __r.

// isResident bool is pointed to by __ir

// Asm args for return values. It's a 4-element vector

#define __ASM_OUT(__t)                                                         \

  ("=" __t(__r.x), "=" __t(__r.y), "=" __t(__r.z), "=" __t(__r.w))

// .. possibly combined with a predicate.

#define __ASM_OUTP(__t) (__L(__ASM_OUT(__t)), "=h"(*__ir))

// Implements a single variant of texture fetch instruction.

#define __IMPL_F1(__rt, __dt, __args, __asm_op, __asm_outs, __asm_args)        \

  template <>                                                                  \

  __device__ __rt __run<__dt>(cudaTextureObject_t __obj, __L(__args)) {        \

    __rt __r;                                                                  \

    asm(__asm_op : __L(__asm_outs) : "l"(__obj), __L(__asm_args));             \

    return __r;                                                                \

  }

// Implements texture fetch instructions for int4/uint4/float4 data types.

#define __IMPL_F3(__args, __asm_op, __ctype, __asm_op_args, __asm_args)        \

  __IMPL_F1(int4, int4, __args, __asm_op ".s32." __ctype "\t" __asm_op_args,   \

            __ASM_OUT("r"), __asm_args)                                        \

  __IMPL_F1(uint4, uint4, __args, __asm_op ".u32." __ctype "\t" __asm_op_args, \

            __ASM_OUT("r"), __asm_args)                                        \

  __IMPL_F1(float4, float4, __args,                                            \

            __asm_op ".f32." __ctype "\t" __asm_op_args, __ASM_OUT("f"),       \

            __asm_args)

// Implements 'sparse' texture fetch instructions for int4/uint4/float4 data

// types. Similar to above, but returns a boolean 'isPresent' value in addition

// to texture data,

#define __IMPL_F3S(__args, __asm_op, __ctype, __asm_op_args, __asm_args)       \

  __IMPL_F1(int4, int4, __args, __asm_op ".s32." __ctype "\t" __asm_op_args,   \

            __ASM_OUTP("r"), __asm_args)                                       \

  __IMPL_F1(uint4, uint4, __args, __asm_op ".u32." __ctype "\t" __asm_op_args, \

            __ASM_OUTP("r"), __asm_args)                                       \

  __IMPL_F1(float4, float4, __args,                                            \

            __asm_op ".f32." __ctype "\t" __asm_op_args, __ASM_OUTP("f"),      \

            __asm_args)

// Similar to F3, but for integer data which is returned as normalized floats.

// Only instantiates fetch functions for int4/uint4.

#define __IMPL_F3N(__args, __asm_op, __ctype, __asm_op_args, __asm_args)       \

  __IMPL_F1(float4, int4, __args, __asm_op ".s32." __ctype "\t" __asm_op_args, \

            __ASM_OUT("r"), __asm_args)                                        \

  __IMPL_F1(float4, uint4, __args,                                             \

            __asm_op ".u32." __ctype "\t" __asm_op_args, __ASM_OUT("r"),       \

            __asm_args)

// Instantiates __tex_fetch_v4 with regular fetch functions.

#define __IMPL_S3I(__op, __args, __asm_op, __ctype, __asm_op_args, __asm_args) \

  template <> struct __tex_fetch_v4<__op> {                                    \

    template <class T>                                                         \

    __device__ static T __run(cudaTextureObject_t __obj, __L(__args));         \

    __IMPL_F3(__args, __asm_op, __ctype, __asm_op_args, __asm_args)            \

  }

// Same, but for sparse ops. Only available on sm_60+

#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 600)

#define __IMPL_S3SI(__op, __args, __asm_op, __ctype, __asm_op_args,            \

                    __asm_args)                                                \

  template <> struct __tex_fetch_v4<__op> {                                    \

    template <class T>                                                         \

    __device__ static T __run(cudaTextureObject_t __obj, __L(__args));         \

    __IMPL_F3S(__args, __asm_op, __ctype, __asm_op_args, __asm_args)           \

  }

#else

#define __IMPL_S3SI(__op, __args, __asm_op, __ctype, __asm_op_args, __asm_args)

#endif

// Same, but for normalized float ops.

#define __IMPL_S3NI(__op, __args, __asm_op, __ctype, __asm_op_args,            \

                    __asm_args)                                                \

  template <> struct __tex_fetch_v4<__op> {                                    \

    template <class T>                                                         \

    __device__ static float4 __run(cudaTextureObject_t __obj, __L(__args));    \

    __IMPL_F3N(__args, __asm_op, __ctype, __asm_op_args, __asm_args)           \

  }

// Regular and normalized float ops share a lot of similarities.  This macro

// instantiates both variants -- normal for __op and normalized for __opn.

#define __IMPL_SI(__op, __opn, __args, __asm_op, __ctype, __asm_op_args,       \

                  __asm_args)                                                  \

  __IMPL_S3I(__op, __args, __asm_op, __ctype, __asm_op_args, __asm_args);      \

  __IMPL_S3NI(__opn, __args, __asm_op, __ctype, __asm_op_args, __asm_args)

// Convenience macros which converts string literal __op into a __Tag,

#define __IMPL_S3(__op, __args, __asm_op, __ctype, __asm_op_args, __asm_args)  \

  __IMPL_S3I(__ID(__op), __args, __asm_op, __ctype, __asm_op_args, __asm_args)

#define __IMPL_S3S(__op, __args, __asm_op, __ctype, __asm_op_args, __asm_args) \

  __IMPL_S3SI(__ID(__op), __args, __asm_op, __ctype, __asm_op_args, __asm_args)

#define __IMPL_S3N(__op, __args, __asm_op, __ctype, __asm_op_args, __asm_args) \

  __IMPL_S3NI(__ID(__op), __args, __asm_op, __ctype, __asm_op_args, __asm_args)

#define __IMPL_S(__op, __opn, __args, __asm_op, __ctype, __asm_op_args,        \

                 __asm_args)                                                   \

  __IMPL_SI(__ID(__op), __ID(__opn), __args, __asm_op, __ctype, __asm_op_args, \

            __asm_args)

// CUDA headers have some 'legacy' texture oprerations that duplicate

// functionality. So, we just inherit it, instead of refining a copy.

#define __IMPL_ALIASI(__op, __opn)                                             \

  template <> struct __tex_fetch_v4<__op> : __tex_fetch_v4<__opn> {}

#define __IMPL_ALIAS(__op, __opn) __IMPL_ALIASI(__ID(__op), __ID(__opn))

// Now we can instantiate everything we need for each specific texture fetch

// variant.

__IMPL_S("__tex1D_v2", "__tex1D_rmnf_v2", (float __x), "tex.1d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5}];", ("f"(__x)));

__IMPL_S("__tex1Dfetch_v2", "__tex1Dfetch_rmnf_v2", (int __x), "tex.1d.v4",

         "s32", "{%0, %1, %2, %3}, [%4, {%5}];", ("r"(__x)));

__IMPL_ALIAS("__itex1D", "__tex1D_v2");

__IMPL_ALIAS("__itex1Dfetch", "__tex1Dfetch_v2");

__IMPL_S("__tex1DGrad_v2", "__tex1DGrad_rmnf_v2",

         (float __x, float __dPdx, float __dPdy), "tex.grad.1d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5}], {%6}, {%7};",

         ("f"(__x), "f"(__dPdx), "f"(__dPdy)));

__IMPL_ALIAS("__itex1DGrad", "__tex1DGrad_v2");

__IMPL_S("__tex1DLayered_v2", "__tex1DLayered_rmnf_v2",

         (float __x, int __layer), "tex.a1d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6}];", ("r"(__layer), "f"(__x)));

__IMPL_ALIAS("__itex1DLayered", "__tex1DLayered_v2");

__IMPL_S("__tex1DLayeredGrad_v2", "__tex1DLayeredGrad_rmnf_v2",

         (float __x, int __layer, float __dPdx, float __dPdy),

         "tex.grad.a1d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6}], {%7}, {%8};",

         ("r"(__layer), "f"(__x), "f"(__dPdx), "f"(__dPdy)));

__IMPL_ALIAS("__itex1DLayeredGrad", "__tex1DLayeredGrad_v2");

__IMPL_S("__tex1DLayeredLod_v2", "__tex1DLayeredLod_rmnf_v2",

         (float __x, int __layer, float __level), "tex.level.a1d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6}], %7;",

         ("r"(__layer), "f"(__x), "f"(__level)));

__IMPL_ALIAS("__itex1DLayeredLod", "__tex1DLayeredLod_v2");

__IMPL_S("__tex1DLod_v2", "__tex1DLod_rmnf_v2", (float __x, float __level),

         "tex.level.1d.v4", "f32", "{%0, %1, %2, %3}, [%4, {%5}], %6;",

         ("f"(__x), "f"(__level)));

__IMPL_ALIAS("__itex1DLod", "__tex1DLod_v2");

// 2D

__IMPL_S("__tex2D_v2", "__tex2D_rmnf_v2", (float __x, float __y), "tex.2d.v4",

         "f32", "{%0, %1, %2, %3}, [%4, {%5, %6}];", ("f"(__x), "f"(__y)));

__IMPL_ALIAS("__itex2D", "__tex2D_v2");

__IMPL_S3S("__itex2D_sparse", (float __x, float __y, unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.2d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7}];\n\t"

           " selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y)));

__IMPL_S("__tex2DGrad_v2", "__tex2DGrad_rmnf_v2",

         (float __x, float __y, const float2 *__dPdx, const float2 *__dPdy),

         "tex.grad.2d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6}], {%7, %8}, {%9, %10};",

         ("f"(__x), "f"(__y), "f"(__dPdx->x), "f"(__dPdx->y), "f"(__dPdy->x),

          "f"(__dPdy->y)));

__IMPL_ALIAS("__itex2DGrad_v2", "__tex2DGrad_v2");

__IMPL_S3S("__itex2DGrad_sparse",

           (float __x, float __y, const float2 *__dPdx, const float2 *__dPdy,

            unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.grad.2d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7}], {%8, %9}, {%10, %11};\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y), "f"(__dPdx->x), "f"(__dPdx->y), "f"(__dPdy->x),

            "f"(__dPdy->y)));

__IMPL_S("__tex2DLayered_v2", "__tex2DLayered_rmnf_v2",

         (float __x, float __y, int __layer), "tex.a2d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}];",

         ("r"(__layer), "f"(__x), "f"(__y)));

__IMPL_ALIAS("__itex2DLayered", "__tex2DLayered_v2");

__IMPL_S3S("__itex2DLayered_sparse",

           (float __x, float __y, int __layer, unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.a2d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}];\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("r"(__layer), "f"(__x), "f"(__y)));

__IMPL_S("__tex2DLayeredGrad_v2", "__tex2DLayeredGrad_rmnf_v2",

         (float __x, float __y, int __layer, const float2 *__dPdx,

          const float2 *__dPdy),

         "tex.grad.a2d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}], {%8, %9}, {%10, %11};",

         ("r"(__layer), "f"(__x), "f"(__y), "f"(__dPdx->x), "f"(__dPdx->y),

          "f"(__dPdy->x), "f"(__dPdy->y)));

__IMPL_ALIAS("__itex2DLayeredGrad_v2", "__tex2DLayeredGrad_v2");

__IMPL_S3S(

    "__itex2DLayeredGrad_sparse",

    (float __x, float __y, int __layer, const float2 *__dPdx,

     const float2 *__dPdy, unsigned char *__ir),

    "{.reg .pred %%p0;\n\t"

    "tex.grad.a2d.v4",

    "f32",

    "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}], {%9, %10}, {%11, %12};\n\t"

    "selp.u16 %4, 1, 0, %%p0; }",

    ("r"(__layer), "f"(__x), "f"(__y), "f"(__dPdx->x), "f"(__dPdx->y),

     "f"(__dPdy->x), "f"(__dPdy->y)));

__IMPL_S("__tex2DLayeredLod_v2", "__tex2DLayeredLod_rmnf_v2",

         (float __x, float __y, int __layer, float __level), "tex.level.a2d.v4",

         "f32", "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}], %8;",

         ("r"(__layer), "f"(__x), "f"(__y), "f"(__level)));

__IMPL_ALIAS("__itex2DLayeredLod", "__tex2DLayeredLod_v2");

__IMPL_S3S("__itex2DLayeredLod_sparse",

           (float __x, float __y, int __layer, float __level,

            unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.level.a2d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}], %9;\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("r"(__layer), "f"(__x), "f"(__y), "f"(__level)));

__IMPL_S("__tex2DLod_v2", "__tex2DLod_rmnf_v2",

         (float __x, float __y, float __level), "tex.level.2d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6}], %7;",

         ("f"(__x), "f"(__y), "f"(__level)));

__IMPL_ALIAS("__itex2DLod", "__tex2DLod_v2");

__IMPL_S3S("__itex2DLod_sparse",

           (float __x, float __y, float __level, unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.level.2d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7}], %8;\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y), "f"(__level)));

// 2D gather is special. Unlike other variants that translate into exactly one

// asm instruction, it uses one of the four different instructions selected by

// __comp.  We implement each instruction variant separately, and dispatch the

// right one from the manually implemented 'umbrella' fetch.

#define __IMPL_2DGATHER(variant, instr)                                        \

  __IMPL_SI(__IDV("__tex2Dgather_v2", variant),                                \

            __IDV("__tex2Dgather_rmnf_v2", variant),                           \

            (float __x, float __y, int __comp), instr, "f32",                  \

            "{%0, %1, %2, %3}, [%4, {%5, %6}];", ("f"(__x), "f"(__y)));        \

  __IMPL_ALIASI(__IDV("__itex2Dgather", variant),                              \

                __IDV("__tex2Dgather_v2", variant));                           \

  __IMPL_S3SI(__IDV("__itex2Dgather_sparse", variant),                         \

              (float __x, float __y, unsigned char *__ir, int __comp),         \

              "{.reg .pred %%p0;\n\t" instr, "f32",                            \

              "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7}];\n\t"                     \

              "selp.u16 %4, 1, 0, %%p0; }",                                    \

              ("f"(__x), "f"(__y)));

__IMPL_2DGATHER(0, "tld4.r.2d.v4");

__IMPL_2DGATHER(1, "tld4.g.2d.v4");

__IMPL_2DGATHER(2, "tld4.b.2d.v4");

__IMPL_2DGATHER(3, "tld4.a.2d.v4");

// Umbrella dispatcher -- calls into specific 2Dgather variant.

template <> struct __tex_fetch_v4<__ID("__tex2Dgather_v2")> {

  template <class __T>

  __device__ static __T __run(cudaTextureObject_t __obj, float __x, float __y,

                              int __comp) {

    switch (__comp) {

    case 0:

      return __tex_fetch_v4<__IDV("__tex2Dgather_v2", 0)>::__run<__T>(

          __obj, __x, __y, __comp);

    case 1:

      return __tex_fetch_v4<__IDV("__tex2Dgather_v2", 1)>::__run<__T>(

          __obj, __x, __y, __comp);

    case 2:

      return __tex_fetch_v4<__IDV("__tex2Dgather_v2", 2)>::__run<__T>(

          __obj, __x, __y, __comp);

    case 3:

      return __tex_fetch_v4<__IDV("__tex2Dgather_v2", 3)>::__run<__T>(

          __obj, __x, __y, __comp);

    }

  }

};

__IMPL_ALIAS("__itex2Dgather", "__tex2Dgather_v2");

template <> struct __tex_fetch_v4<__ID("__tex2Dgather_rmnf_v2")> {

  template <class __T>

  __device__ static float4 __run(cudaTextureObject_t __obj, float __x,

                                 float __y, int __comp) {

    switch (__comp) {

    case 0:

      return __tex_fetch_v4<__IDV("__tex2Dgather_rmnf_v2", 0)>::__run<__T>(

          __obj, __x, __y, __comp);

    case 1:

      return __tex_fetch_v4<__IDV("__tex2Dgather_rmnf_v2", 1)>::__run<__T>(

          __obj, __x, __y, __comp);

    case 2:

      return __tex_fetch_v4<__IDV("__tex2Dgather_rmnf_v2", 2)>::__run<__T>(

          __obj, __x, __y, __comp);

    case 3:

      return __tex_fetch_v4<__IDV("__tex2Dgather_rmnf_v2", 3)>::__run<__T>(

          __obj, __x, __y, __comp);

    }

  }

};

#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 600)

template <> struct __tex_fetch_v4<__ID("__itex2Dgather_sparse")> {

  template <class __T>

  __device__ static __T __run(cudaTextureObject_t __obj, float __x, float __y,

                              unsigned char *__ir, int __comp) {

    switch (__comp) {

    case 0:

      return __tex_fetch_v4<__IDV("__itex2Dgather_sparse", 0)>::__run<__T>(

          __obj, __x, __y, __ir, __comp);

    case 1:

      return __tex_fetch_v4<__IDV("__itex2Dgather_sparse", 1)>::__run<__T>(

          __obj, __x, __y, __ir, __comp);

    case 2:

      return __tex_fetch_v4<__IDV("__itex2Dgather_sparse", 2)>::__run<__T>(

          __obj, __x, __y, __ir, __comp);

    case 3:

      return __tex_fetch_v4<__IDV("__itex2Dgather_sparse", 3)>::__run<__T>(

          __obj, __x, __y, __ir, __comp);

    }

  }

};

#endif

// 3D

__IMPL_S("__tex3D_v2", "__tex3D_rmnf_v2", (float __x, float __y, float __z),

         "tex.3d.v4", "f32", "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}];",

         ("f"(__x), "f"(__y), "f"(__z)));

__IMPL_ALIAS("__itex3D", "__tex3D_v2");

__IMPL_S3S("__itex3D_sparse",

           (float __x, float __y, float __z, unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.3d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}];\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y), "f"(__z)));

__IMPL_S("__tex3DGrad_v2", "__tex3DGrad_rmnf_v2",

         (float __x, float __y, float __z, const float4 *__dPdx,

          const float4 *__dPdy),

         "tex.grad.3d.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}], "

         "{%8, %9, %10, %10}, {%11, %12, %13, %13};",

         ("f"(__x), "f"(__y), "f"(__z), "f"(__dPdx->x), "f"(__dPdx->y),

          "f"(__dPdx->z), "f"(__dPdy->x), "f"(__dPdy->y), "f"(__dPdy->z)));

__IMPL_ALIAS("__itex3DGrad_v2", "__tex3DGrad_v2");

__IMPL_S3S("__itex3DGrad_sparse",

           (float __x, float __y, float __z, const float4 *__dPdx,

            const float4 *__dPdy, unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.grad.3d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}], "

           "{%9, %10, %11, %11}, {%12, %13, %14, %14};\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y), "f"(__z), "f"(__dPdx->x), "f"(__dPdx->y),

            "f"(__dPdx->z), "f"(__dPdy->x), "f"(__dPdy->y), "f"(__dPdy->z)));

__IMPL_S("__tex3DLod_v2", "__tex3DLod_rmnf_v2",

         (float __x, float __y, float __z, float __level), "tex.level.3d.v4",

         "f32", "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}], %8;",

         ("f"(__x), "f"(__y), "f"(__z), "f"(__level)));

__IMPL_ALIAS("__itex3DLod", "__tex3DLod_v2");

__IMPL_S3S("__itex3DLod_sparse",

           (float __x, float __y, float __z, float __level,

            unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.level.3d.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}], %9;\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y), "f"(__z), "f"(__level)));

// Cubemap

__IMPL_S("__texCubemap_v2", "__texCubemap_rmnf_v2",

         (float __x, float __y, float __z), "tex.cube.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}];",

         ("f"(__x), "f"(__y), "f"(__z)));

__IMPL_ALIAS("__itexCubemap", "__texCubemap_v2");

__IMPL_S3S("__itexCubemap_sparse",

           (float __x, float __y, float __z, unsigned char *__ir),

           "{.reg .pred %%p0;\n\t"

           "tex.cube.v4",

           "f32",

           "{%0, %1, %2, %3}|%%p0, [%5, {%6, %7, %8, %8}];\n\t"

           "selp.u16 %4, 1, 0, %%p0; }",

           ("f"(__x), "f"(__y), "f"(__z)));

__IMPL_S("__texCubemapGrad_v2", "__texCubemapGrad_rmnf_v2",

         (float __x, float __y, float __z, const float4 *__dPdx,

          const float4 *__dPdy),

         "tex.grad.cube.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}], "

         "{%8, %9, %10, %10}, {%11, %12, %13, %13};",

         ("f"(__x), "f"(__y), "f"(__z), "f"(__dPdx->x), "f"(__dPdx->y),

          "f"(__dPdx->z), "f"(__dPdy->x), "f"(__dPdy->y), "f"(__dPdy->z)));

__IMPL_ALIAS("__itexCubemapGrad_v2", "__texCubemapGrad_v2");

__IMPL_S("__texCubemapLayered_v2", "__texCubemapLayered_rmnf_v2",

         (float __x, float __y, float __z, int __layer), "tex.acube.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %8}];",

         ("r"(__layer), "f"(__x), "f"(__y), "f"(__z)));

__IMPL_ALIAS("__itexCubemapLayered", "__texCubemapLayered_v2");

__IMPL_S("__texCubemapLayeredGrad_v2", "__texCubemapLayeredGrad_rmnf_v2",

         (float __x, float __y, float __z, int __layer, const float4 *__dPdx,

          const float4 *__dPdy),

         "tex.grad.acube.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %8}], "

         "{%9, %10, %11, %11}, {%12, %13, %14, %14};",

         ("r"(__layer), "f"(__x), "f"(__y), "f"(__z), "f"(__dPdx->x),

          "f"(__dPdx->y), "f"(__dPdx->z), "f"(__dPdy->x), "f"(__dPdy->y),

          "f"(__dPdy->z)));

__IMPL_ALIAS("__itexCubemapLayeredGrad_v2", "__texCubemapLayeredGrad_v2");

__IMPL_S("__texCubemapLayeredLod_v2", "__texCubemapLayeredLod_rmnf_v2",

         (float __x, float __y, float __z, int __layer, float __level),

         "tex.level.acube.v4", "f32",

         "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %8}], %9;",

         ("r"(__layer), "f"(__x), "f"(__y), "f"(__z), "f"(__level)));

__IMPL_ALIAS("__itexCubemapLayeredLod", "__texCubemapLayeredLod_v2");

__IMPL_S("__texCubemapLod_v2", "__texCubemapLod_rmnf_v2",

         (float __x, float __y, float __z, float __level), "tex.level.cube.v4",

         "f32", "{%0, %1, %2, %3}, [%4, {%5, %6, %7, %7}], %8;",

         ("f"(__x), "f"(__y), "f"(__z), "f"(__level)));

__IMPL_ALIAS("__itexCubemapLod", "__texCubemapLod_v2");

// Helper class for extracting slice of data from V4 fetch results.

template <class __DestT, class __SrcT> struct __convert {

  template <int __NElements = sizeof(__DestT) /

                              sizeof(typename __TypeInfoT<__DestT>::__base_t)>

  __device__ static __DestT __run(__SrcT __v);

  template <> __device__ static __DestT __run<1>(__SrcT __v) { return {__v.x}; }

  template <> __device__ static __DestT __run<2>(__SrcT __v) {

    return {__v.x, __v.y};

  }

  template <> __device__ static __DestT __run<3>(__SrcT __v) {

    return {__v.x, __v.y, __v.z};

  }

  template <> __device__ static __DestT __run<4>(__SrcT __v) {

    return {__v.x, __v.y, __v.z, __v.w};

  }

};

// These are the top-level function overloads the __nv_tex_surf_handler expands

// to.  Each overload deals with one of the several ways __nv_tex_surf_handler

// is called by CUDA headers. In the end, each of the overloads does the same

// job -- it figures out which `__tex_fetch_v4::run` variant should be used to

// fetch texture data and which `__convert::run` is needed to convert it into

// appropriate return type.

// __nv_tex_surf_handler("__tex...", &ret, cudaTextureObject_t handle, args...);

//   Data type and return type are based on ret.

template <class __op, class __T, class... __Args>

__device__ static void __tex_fetch(__T *__ptr, cudaTextureObject_t __handle,

                                   __Args... __args) {

  using __FetchT = typename __TypeInfoT<__T>::__fetch_t;

  *__ptr = __convert<__T, __FetchT>::__run(

      __tex_fetch_v4<__op>::template __run<__FetchT>(__handle, __args...));

}

#if CUDA_VERSION < 12000

// texture<> objects get magically converted into a texture reference.  However,

// there's no way to convert them to cudaTextureObject_t on C++ level. So, we

// cheat a bit and use inline assembly to do it. It costs us an extra register

// and a move, but that is easy for ptxas to optimize away.

template <class __T>

__device__ cudaTextureObject_t __tex_handle_to_obj(__T __handle) {

  cudaTextureObject_t __obj;

  asm("mov.b64 %0, %1; " : "=l"(__obj) : "l"(__handle));

  return __obj;

}

// __nv_tex_surf_handler ("__tex...", &ret, textureReference, args...);

//   Data type and return type is based on ret.

template <class __op, class __T, class __HandleT, class... __Args>

__device__ static void __tex_fetch(__T *__ptr, __HandleT __handle,

                                   __Args... __args) {

  using __FetchT = typename __TypeInfoT<__T>::__fetch_t;

  *__ptr = __convert<__T, __FetchT>::__run(

      __tex_fetch_v4<__op>::template __run<__FetchT>(

          __tex_handle_to_obj(__handle), __args...));

}

// __nv_tex_surf_handler ("__tex...", &type_dummy, &ret, texture<...>, args...);

// cudaReadModeNormalizedFloat fetches always return float4.

template <class __op, class __DataT, class __RetT, int __TexT, class... __Args>

__device__ static void

__tex_fetch(__DataT *, __RetT *__ptr,

            texture<__DataT, __TexT, cudaReadModeNormalizedFloat> __handle,

            __Args... __args) {

  using __FetchT = typename __TypeInfoT<__DataT>::__fetch_t;

  *__ptr = __convert<__RetT, float4>::__run(

      __tex_fetch_v4<__op>::template __run<__FetchT>(

          __tex_handle_to_obj(__handle), __args...));

}

// __nv_tex_surf_handler ("__tex...", &type_dummy, &ret, texture<...>, args...);

// For cudaReadModeElementType fetch return type is based on type_dummy.

template <class __op, class __DataT, class __RetT, int __TexT, class... __Args>

__device__ static void

__tex_fetch(__DataT *, __RetT *__ptr,

            texture<__DataT, __TexT, cudaReadModeElementType> __handle,

            __Args... __args) {

  using __FetchT = typename __TypeInfoT<__DataT>::__fetch_t;

  *__ptr = __convert<__RetT, __FetchT>::__run(

      __tex_fetch_v4<__op>::template __run<__FetchT>(

          __tex_handle_to_obj(__handle), __args...));

}

#endif // CUDA_VERSION

} // namespace __cuda_tex

} // namespace

#pragma pop_macro("__ASM_OUT")

#pragma pop_macro("__ASM_OUTP")

#pragma pop_macro("__Args")

#pragma pop_macro("__ID")

#pragma pop_macro("__IDV")

#pragma pop_macro("__IMPL_2DGATHER")

#pragma pop_macro("__IMPL_ALIAS")

#pragma pop_macro("__IMPL_ALIASI")

#pragma pop_macro("__IMPL_F1")

#pragma pop_macro("__IMPL_F3")

#pragma pop_macro("__IMPL_F3N")

#pragma pop_macro("__IMPL_F3S")

#pragma pop_macro("__IMPL_S")

#pragma pop_macro("__IMPL_S3")

#pragma pop_macro("__IMPL_S3I")

#pragma pop_macro("__IMPL_S3N")

#pragma pop_macro("__IMPL_S3NI")

#pragma pop_macro("__IMPL_S3S")

#pragma pop_macro("__IMPL_S3SI")

#pragma pop_macro("__IMPL_SI")

#pragma pop_macro("__L")