5.3. C++ 语言支持（C++ Language Support）

    __host__ __device__ clock_t   clock();
    __device__          long long clock64();

    __host__ __device__ __tile__ int printf(const char* format[, arg, ...]);

    #include <stdio.h>

    __global__ void helloCUDA(float value) {
        printf("Hello thread %d, value=%f\n", threadIdx.x, value);
    }

    int main() {
        helloCUDA<<<1, 5>>>(1.2345f);
        cudaDeviceSynchronize();
        return 0;
    }

    #include <stdio.h>

    __global__ void helloCUDA(float value) {
        if (threadIdx.x == 0)
            printf("Hello thread %d, value=%f\n", threadIdx.x, value);
    }

    int main() {
        helloCUDA<<<1, 5>>>(1.2345f);
        cudaDeviceSynchronize();
        return 0;
    }

    #include "cuda_tile.h"
    #include <cstdio>

    namespace ct = cuda::tiles;

    __tile_global__ void kernel() {
      auto ints = ct::iota<ct::tile<int, ct::shape<4, 4>>>();
      printf("%i\n", ints);
    }

    int main() {
      kernel<<<1,1>>>();
      cudaDeviceSynchronize();
      return 0;
    }

    __host__ __device__ __tile__ void* memcpy(void* dest, const void* src, size_t size);

    __host__ __device__ __tile__ void* memset(void* ptr, int value, size_t size);

    __host__ __device__ void* malloc(size_t size);
    // or cuda::std::malloc(), cuda::std::calloc() in the <cuda/std/cstdlib> header

    __device__ void* __nv_aligned_device_malloc(size_t size, size_t align);
    // or cuda::std::aligned_alloc() in the <cuda/std/cstdlib> header

    __host__ __device__ void free(void* ptr);
    // or cuda::std::free() in the <cuda/std/cstdlib> header

    #include <stdlib.h>
    #include <stdio.h>

    __global__ void single_thread_allocation_kernel() {
        size_t size = 123;
        char*  ptr  = (char*) malloc(size);
        memset(ptr, 0, size);
        printf("Thread %d got pointer: %p\n", threadIdx.x, ptr);
        free(ptr);
    }

    int main() {
        // Set a heap size of 128 megabytes.
        // Note that this must be done before any kernel is launched.
        cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024);
        single_thread_allocation_kernel<<<1, 5>>>();
        cudaDeviceSynchronize();
        return 0;
    }

    Thread 0 got pointer: 0x20d5ffe20
    Thread 1 got pointer: 0x20d5ffec0
    Thread 2 got pointer: 0x20d5fff60
    Thread 3 got pointer: 0x20d5f97c0
    Thread 4 got pointer: 0x20d5f9720

    #include <stdlib.h>

    __global__ void block_level_allocation_kernel() {
        __shared__ int* data;
        // The first thread in the block performs the allocation and shares the pointer
        // with all other threads through shared memory, so that access can be coalesced.
        if (threadIdx.x == 0) {
            size_t size = blockDim.x * 64; // 64 bytes per thread are allocated.
            data = (int*) malloc(size);
        }
        __syncthreads();
        // Check for failure
        if (data == nullptr)
            return;

        // Threads index into the memory, ensuring coalescence
        for (int i = 0; i < 64; ++i)
            data[i * blockDim.x + threadIdx.x] = threadIdx.x;
        // Ensure all threads complete before freeing
        __syncthreads();

        // Only one thread may free the memory!
        if (threadIdx.x == 0)
            free(data);
    }

    int main() {
        cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024);
        block_level_allocation_kernel<<<10, 128>>>();
        cudaDeviceSynchronize();
        return 0;
    }

    #include <stdlib.h>
    #include <stdio.h>

    const int NUM_BLOCKS = 20;

    __device__ int* data_ptrs[NUM_BLOCKS]; // Per-block pointer

    __global__ void allocate_memory_kernel() {
        // Only the first thread in the block performs the allocation
        // since we need only one allocation per block.
        if (threadIdx.x == 0)
            data_ptrs[blockIdx.x] = (int*) malloc(blockDim.x * 4);
        __syncthreads();
        // Check for failure
        if (data_ptrs[blockIdx.x] == nullptr)
            return;
        // Zero the data with all threads in parallel
        data_ptrs[blockIdx.x][threadIdx.x] = 0;
    }

    // Simple example: store the thread ID into each element
    __global__ void use_memory_kernel() {
        int* ptr = data_ptrs[blockIdx.x];
        if (ptr != nullptr)
            ptr[threadIdx.x] += threadIdx.x;
    }

    // Print the content of the buffer before freeing it
    __global__ void free_memory_kernel() {
        int* ptr = data_ptrs[blockIdx.x];
        if (ptr != nullptr)
            printf("Block %d, Thread %d: final value = %d\n",
                blockIdx.x, threadIdx.x, ptr[threadIdx.x]);
        // Only free from one thread!
        if (threadIdx.x == 0)
            free(ptr);
    }

    int main() {
        cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
        // Allocate memory
        allocate_memory_kernel<<<NUM_BLOCKS, 10>>>();

        // Use memory
        use_memory_kernel<<<NUM_BLOCKS, 10>>>();
        use_memory_kernel<<<NUM_BLOCKS, 10>>>();
        use_memory_kernel<<<NUM_BLOCKS, 10>>>();

        // Free memory
        free_memory_kernel<<<NUM_BLOCKS, 10>>>();
        cudaDeviceSynchronize();
        return 0;
    }

    __host__ __device__ void* alloca(size_t size);

    __device__ void device_function(int num_items) {
        int4* ptr = (int4*) alloca(num_items * sizeof(int4));
        // use of ptr
        ...
    }

    auto global_lambda = [](){ return 0; }; // __host__

    void host_function() {
      auto lambda1 = [](){ return 1; };   // __host__
      [](){ return 3; };                  // __host__, closure type (body of a lambda expression)
    }

    __device__ void device_function() {
      auto lambda2 = [](){ return 2; };   // __device__
    }

    __global__ void kernel_function(void) {
      auto lambda3 = [](){ return 3; };   // __device__
    }

    __host__ __device__ void host_device_function() {
      auto lambda4 = [](){ return 4; };   // __host__ __device__
    }

    __tile__ void tile_function() {
      auto lambda5 = [](){ return 5; };   // __tile__
    }

    __tile__ __device__ void tile_device_function() {
      auto lambda6 = [](){ return 6; };   // __tile__ __device__
    }

    using function_ptr_t = int (*)();

    __device__ void device_function(float          value,
                                    function_ptr_t ptr = [](){ return 4; } /* __host__ */) {}

    template <typename T>
     __global__ void kernel(T input) {}

     __device__ void device_function() {
         // device kernel call requires separate compilation (-rdc=true flag)
         kernel<<<1, 1>>>([](){});
         kernel<<<1, 1>>>([] __device__() {});          // extended lambda
         kernel<<<1, 1>>>([] __host__ __device__() {}); // extended lambda
     }

     auto global_lambda = [] __host__ __device__() {};

     void host_function() {
         kernel<<<1, 1>>>([] __device__() {});          // CORRECT, extended lambda
         kernel<<<1, 1>>>([] __host__ __device__() {}); // CORRECT, extended lambda
     //  kernel<<<1, 1>>>([](){});                      // ERROR, closure type with host execution space
     //  kernel<<<1, 1>>>(global_lambda);               // ERROR, extended lambda, but at global scope
     }

    void host_function() {
        auto lambda1 = [] {};                      // NOT an extended lambda: no explicit execution space annotations
        auto lambda2 = [] __device__ {};           // extended lambda
        auto lambda3 = [] __host__ __device__ {};  // extended lambda
        auto lambda4 = [] __host__ {};             // NOT an extended lambda
    }

    __host__ __device__ void host_device_function() {
        auto lambda1 = [] {};                      // NOT an extended lambda: no explicit execution space annotations
        auto lambda2 = [] __device__ {};           // extended lambda
        auto lambda3 = [] __host__ __device__ {};  // extended lambda
        auto lambda4 = [] __host__ {};             // NOT an extended lambda
    }

    __device__ void device_function() {
        // none of the lambdas within this function are extended lambdas,
        // because the enclosing function is not a __host__ or __host__ __device__  function.
        auto lambda1 = [] {};
        auto lambda2 = [] __device__ {};
        auto lambda3 = [] __host__ __device__ {};
        auto lambda4 = [] __host__ {};
    }

    auto global_lambda = [] __host__ __device__ { }; // NOT an extended lambda because it is not defined
                                                     // within a __host__ or __host__ __device__ function

    bool __nv_is_extended_device_lambda_closure_type(type);

    bool __nv_is_extended_device_lambda_with_preserved_return_type(type);

    bool __nv_is_extended_host_device_lambda_closure_type(type);

    auto lambda0 = [] __host__ __device__ { };

    void host_function() {
        auto lambda1 = [] { };
        auto lambda2 = [] __device__ { };
        auto lambda3 = [] __host__ __device__ { };
        auto lambda4 = [] __device__ () -> double { return 3.14; }
        auto lambda5 = [] __device__ (int x) -> decltype(&x) { return 0; }

        using lambda0_t = decltype(lambda0);
        using lambda1_t = decltype(lambda1);
        using lambda2_t = decltype(lambda2);
        using lambda3_t = decltype(lambda3);
        using lambda4_t = decltype(lambda4);
        using lambda5_t = decltype(lambda5);

        // 'lambda0' is not an extended lambda because it is defined outside function scope
        static_assert(!__nv_is_extended_device_lambda_closure_type(lambda0_t));
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda0_t));
        static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda0_t));

        // 'lambda1' is not an extended lambda because it has no execution space annotations
        static_assert(!__nv_is_extended_device_lambda_closure_type(lambda1_t));
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda1_t));
        static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda1_t));

        // 'lambda2' is an extended device-only lambda
        static_assert(__nv_is_extended_device_lambda_closure_type(lambda2_t));
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda2_t));
        static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda2_t));

        // 'lambda3' is an extended host-device lambda
        static_assert(!__nv_is_extended_device_lambda_closure_type(lambda3_t));
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda3_t));
        static_assert(__nv_is_extended_host_device_lambda_closure_type(lambda3_t));

        // 'lambda4' is an extended device-only lambda with preserved return type
        static_assert(__nv_is_extended_device_lambda_closure_type(lambda4_t));
        static_assert(__nv_is_extended_device_lambda_with_preserved_return_type(lambda4_t));
        static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda4_t));

        // 'lambda5' is not an extended device-only lambda with preserved return type
        // because it references the operator()'s parameter types in the trailing return type.
        static_assert(__nv_is_extended_device_lambda_closure_type(lambda5_t));
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda5_t));
        static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda5_t));
    }

    void host_function() {
        auto lambda1 = [] __device__ { }; // enclosing function for lambda1 is "host_function()"
        auto lambda2 = [] {
            auto lambda3 = [] {
                auto lambda4 = [] __host__ __device__ { }; // enclosing function for lambda4 is "host_function"
            };
        };
    }

    auto global_lambda = [] {
        auto lambda5 = [] __host__ __device__ { }; // enclosing function for lambda5 does not exist
    };

    void host_function() {
        auto lambda1 = [] __host__ __device__  {
             // ERROR, extended lambda defined within another extended lambda
            auto lambda2 = [] __host__ __device__ { };
        };
    }

    void host_function() {
        auto lambda1 = [] (auto) {
             // ERROR, extended lambda defined within a generic lambda
            auto lambda2 = [] __host__ __device__ { };
        };
    }

    auto lambda1 = []  {
        // ERROR, outer enclosing lambda is not defined within a non-lambda-operator() function
        auto lambda2 = [] __host__ __device__ { };
    };

    void host_function() {
        auto lambda1 = [] __device__ { return 0; }; // OK
        {
            auto lambda2 = [] __device__          { return 0; }; // OK
            auto lambda3 = [] __device__ __host__ { return 0; }; // OK
        }
    }

    struct MyStruct1 {
        MyStruct1() {
            auto lambda4 = [] __device__ { return 0; }; // ERROR, address of the enclosing function is not accessible
        }
    };

    class MyStruct2 {
        void foo() {
            auto temp1 = [] __device__ { return 10; }; // ERROR, enclosing function has private access in parent class
        }

        struct MyStruct3 {
            void foo() {
                auto temp1 = [] __device__ { return 10; };  // ERROR, enclosing class MyStruct3 has private access in its parent class
            }
        };
    };

    template <typename T>
    struct A {
        using Bar = void;
        void test();
    };

    template<>
    struct A<void> { };

    template <typename Bar>
    void A<Bar>::test() {
        // In code sent to host compiler, nvcc will inject an address expression here, of the form:
        //   (void (A< Bar> ::*)(void))(&A::test))
        //  However, the class typedef 'Bar' (to void) shadows the template argument 'Bar',
        //  causing the address expression in A<int>::test to actually refer to:
        //    (void (A< void> ::*)(void))(&A::test))
        //  which doesn't take the address of the enclosing routine 'A<int>::test' correctly.
        auto lambda1 = [] __host__ __device__ { return 4; };
    }

    int main() {
        A<int> var;
        var.test();
    }

    void host_function() {
        struct MyStruct {
            void bar() {
                // ERROR, bar() is member of a class that is local to a function
                auto lambda2 = [] __host__ __device__ { return 0; };
            }
        };
    }

    auto host_function() {
        // ERROR, the return type of host_function() is deduced
        auto lambda3 = [] __host__ __device__ { return 0; };
    }

    void host_function() {
        // ERROR, __host__ __device__ extended lambdas cannot be a generic lambda
        auto lambda1 = [] __host__ __device__ (auto i) { return i; };

        // ERROR, a host-device extended lambda cannot be a generic lambda
        auto lambda2 = [] __host__ __device__ (auto... i) {
            return sizeof...(i);
        };
    }

    template <template <typename...> class T,
              typename... P1,
              typename... P2>
    void bar1(const T<P1...>, const T<P2...>) {
        // ERROR, enclosing function has multiple parameter packs
        auto lambda = [] __device__ { return 10; };
    }

    template <template <typename...> class T,
              typename... P1,
              typename    T2>
    void bar2(const T<P1...>, T2) {
        // ERROR, for enclosing function, the parameter pack is not last in the template parameter list
        auto lambda = [] __device__ { return 10; };
    }

    template <typename T, T>
    void bar3() {
        // ERROR, for enclosing function, the second template parameter is not named
        auto lambda = [] __device__ { return 10; };
    }

    template <typename T>
    void bar4() {
        auto lambda1 = [] __device__ { return 10; };
    }

    class MyStruct {
        struct MyNestedStruct {};

        friend int main();
    };

    int main() {
        struct MyLocalStruct {};
        // ERROR, enclosing function for device lambda in bar4() is instantiated with a type local to main
        bar4<MyLocalStruct>();

        // ERROR, enclosing function for device lambda in bar4 is instantiated with a type
        //        that is a private member of a class
        bar4<MyStruct::MyNestedStruct>();
    }

    void host_function() {
        // CORRECT, an init-capture is allowed for an extended device-only lambda
        auto lambda1 = [x = 1] __device__ () { return x; };

        // ERROR, an init-capture is not allowed for an extended host-device lambda
        auto lambda2 = [x = 1] __host__ __device__ () { return x; };

        int a = 1;
        // ERROR, an extended __device__ lambda cannot capture variables by reference
        auto lambda3 = [&a] __device__ () { return a; };

        // ERROR, by-reference capture is not allowed for an extended device-only lambda
        auto lambda4 = [&x = a] __device__ () { return x; };

        struct MyStruct {};
        MyStruct s1;
        // ERROR, a type local to a function cannot be used in the type of a captured variable
        auto lambda6 = [s1] __device__ () { };

        // ERROR, an init-capture cannot be of type std::initializer_list
        auto lambda7 = [x = {11}] __device__ () { };

        std::initializer_list<int> b = {11,22,33};
        // ERROR, an init-capture cannot be of type std::initializer_list
        auto lambda8 = [x = b] __device__ () { };

        int  var     = 4;
        auto lambda9 = [=] __device__ {
            int result = 0;
            if constexpr(false) {
                //ERROR, An extended device-only lambda cannot first-capture 'var' in if-constexpr context
                result += var;
            }
            return result;
        };

        auto lambda10 = [var] __device__ {
            int result = 0;
            if constexpr(false) {
                // CORRECT, 'var' already listed in explicit capture list for the extended lambda
                result += var;
            }
            return result;
        };

        auto lambda11 = [=] __device__ {
            int result = var;
            if constexpr(false) {
                // CORRECT, 'var' already implicit captured outside the 'if-constexpr' block
                result += var;
            }
            return result;
        };
    }

    template <typename T>
    __global__ void kernel(T in) { in(); }

    __host__ __device__ void host_device_function() {
        // ERROR, the number and relative declaration order of
        //        extended lambdas depend on __CUDA_ARCH__
    #if defined(__CUDA_ARCH__)
        auto lambda1 = [] __device__ { return 0; };
        auto lambda2 = [] __host__ __device__ { return 10; };
    #endif
        auto lambda3 = [] __device__ { return 4; };
        kernel<<<1, 1>>>(lambda3);
    }

    #include <cuda/std/type_traits>

    const char& getRef(const char* p) { return *p; }

    void foo() {
        auto lambda1 = [] __device__ { return "10"; };

        // ERROR, attempt to extract the return type of a device lambda in host code
        cuda::std::result_of<decltype(lambda1)()>::type xx1 = "abc";

        auto lambda2 = [] __host__ __device__ { return "10"; };

        // CORRECT, lambda2 represents a host-device extended lambda
        cuda::std::result_of<decltype(lambda2)()>::type xx2 = "abc";

        auto lambda3 = [] __device__ () -> const char* { return "10"; };

        // CORRECT, lambda3 represents a device extended lambda with preserved return type
        cuda::std::result_of<decltype(lambda3)()>::type xx2 = "abc";
        static_assert(cuda::std::is_same_v<cuda::std::result_of<decltype(lambda3)()>::type, const char*>);

        auto lambda4 = [] __device__ (char x) -> decltype(getRef(&x)) { return 0; };
        // lambda4's return type is not preserved because it references the operator()'s
        // parameter types in the trailing return type.
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(decltype(lambda4)));
    }

    __device__ int result;

    template <typename T>
    __global__ void kernel(T in) { result = in(); }

    void foo(void) {
        int x1 = 1;
        // ERROR, "x1" is only captured when __CUDA_ARCH__ is defined.
        auto lambda1 = [=] __host__ __device__ {
    #ifdef __CUDA_ARCH__
            return x1 + 1;
    #else
            return 10;
    #endif
        };
        kernel<<<1, 1>>>(lambda1);
    }

    template <typename T>
    __global__ void kernel(T in) {
        int (*fp)(double) = in;
        fp(0); // CORRECT, conversion in device code is supported
        auto lambda1 = [](double) { return 1; };
    }

    void foo() {
        auto lambda_device      = [] __device__ (double) { return 1; };
        auto lambda_host_device = [] __host__ __device__ (double) { return 1; };
        kernel<<<1, 1>>>(lambda_device);
        kernel<<<1, 1>>>(lambda_host_device);

        // CORRECT, conversion for a __host__ __device__ lambda is supported in host code
        int (*fp1)(double) = lambda_host_device;

        // ERROR, conversion for a device lambda is not supported in host code
        int (*fp2)(double) = lambda_device;
    }

    #include <cstdio>
    #include <type_traits>

    template <bool b>
    void __global__ kernel() { printf("hi"); }

    template <typename T>
    void kernel_launch() {
        // ERROR, this kernel launch may fail, because CUDA frontend compiler and host compiler
        //        may disagree on the result of std::is_trivially_copyable_v trait on the
        //        closure type of the extended lambda
        kernel<std::is_trivially_copyable_v<T>><<<1,1>>>();
        cudaDeviceSynchronize();
    }

    int main() {
        int  x       = 0;
        auto lambda1 = [=] __host__ __device__ () { return x; };
        kernel_launch<decltype(lambda1)>();
    }

    #include <cstdio>

    template <typename T>
    __global__ void foo(T in) { printf("value = %d\n", in()); }

    struct MyStruct {
        int var;

        __host__ __device__ MyStruct() : var(10) {};

        void run() {
            auto lambda1 = [=] __device__ {
                // reference to "var" causes the 'this' pointer (MyStruct*) to be captured by value
                return var + 1;
            };
            // Kernel launch fails at run time because 'this->var' is not accessible from the GPU
            foo<<<1, 1>>>(lambda1);
            cudaDeviceSynchronize();
        }
    };

    int main() {
        MyStruct s1;
        s1.run();
    }

    #include <cstdio>

    template <typename T>
    __global__ void foo(T in) { printf("\n value = %d", in()); }

    struct MyStruct {
        int var;
        __host__ __device__ MyStruct() : var(10) { };

        void run() {
            // note the "*this" capture specification
            auto lambda1 = [=, *this] __device__ {
                // reference to "var" causes the object denoted by '*this' to be captured by
                // value, and the GPU code will access 'copy_of_star_this->var'
                return var + 1;
            };
            // Kernel launch succeeds
            foo<<<1, 1>>>(lambda1);
            cudaDeviceSynchronize();
        }
    };

    int main() {
        MyStruct s1;
        s1.run();
    }

    struct MyStruct {
        int var;
        __host__ __device__ MyStruct() : var(10) { };

        void host_function() {
            // CORRECT, use in an extended device-only lambda
            auto lambda1 = [=, *this] __device__ { return var; };

            // Use in an extended host-device lambda
            // Error if *this capture not enabled by language dialect
            auto lambda2 = [=, *this] __host__ __device__ { return var; };

            // Use in an non-annotated lambda in host function
            // Error if *this capture not enabled by language dialect
            auto lambda3 = [=, *this]  { return var; };
        }

        __device__ void device_function() {
            // CORRECT, use in a lambda defined in a device-only function
            auto lambda1 = [=, *this] __device__ { return var; };

            // CORRECT, use in a lambda defined in a device-only function
            auto lambda2 = [=, *this] __host__ __device__ { return var; };

            // CORRECT, use in a lambda defined in a device-only function
            auto lambda3 = [=, *this]  { return var; };
        }

        __host__ __device__ void host_device_function() {
            // CORRECT, use in an extended device-only lambda
            auto lambda1 = [=, *this] __device__ { return var; };

            // Use in an extended host-device lambda
            // Error if *this capture not enabled by language dialect
            auto lambda2 = [=, *this] __host__ __device__ { return var; };

            // Use in an unannotated lambda in a host-device function
            // Error if *this capture not enabled by language dialect
            auto lambda3 = [=, *this]  { return var; };
        }
    };

    namespace N1 {

    struct MyStruct {};

    template <typename T>
    void my_function(T);

    }; // namespace N1

    namespace N2 {

    template <typename T>
    int my_function(T);

    template <typename T>
    void run(T in) { my_function(in); }

    } // namespace N2

    void bar(N1::MyStruct in) {
        // For extended device-only lambda, the code sent to the host compiler is replaced with
        // the placeholder type instantiation expression
        //    ' __nv_dl_wrapper_t< __nv_dl_tag<void (*)(N1::MyStruct in),(&bar),1> > { }'
        //
        // As a result, the namespace 'N1' participates in ADL lookup of the
        // call to "my_function()" in the body of N2::run, causing ambiguity.
        auto lambda1 = [=] __device__ { };
        N2::run(lambda1);
    }

    #include <nvfunctional>

    __host__            int host_function()        { return 1; }
    __device__          int device_function()      { return 2; }
    __host__ __device__ int host_device_function() { return 3; }

    __global__ void kernel(int* result) {
        nvstd::function<int()> fn1 = device_function;
        nvstd::function<int()> fn2 = host_device_function;
        nvstd::function<int()> fn3 = [](){ return 10; };
        *result                    = fn1() + fn2() + fn3();
    }

    __host__ __device__ void host_device_test(int* result) {
        nvstd::function<int()> fn1 = host_device_function;
        nvstd::function<int()> fn2 = [](){ return 10; };
        *result                    = fn1() + fn2();
    }

    __host__ void host_test(int* result) {
        nvstd::function<int()> fn1 = host_function;
        nvstd::function<int()> fn2 = host_device_function;
        nvstd::function<int()> fn3 = [](){ return 10; };
        *result                    = fn1() + fn2() + fn3();
    }

    #include <nvfunctional>

    __device__ int device_function() { return 1; }
    __host__   int host_function() { return 3; }
    auto       lambda_host  = [] { return 0; };

    __global__ void k() {
        nvstd::function<int()> fn1 = host_function; // ERROR, initialized with address of __host__ function
        nvstd::function<int()> fn2 = lambda_host;   // ERROR, initialized with address of functor with
                                                    //        __host__ operator() function
    }

    __global__ void kernel(nvstd::function<int()> f1) {}

    void foo(void) {
        auto lambda_device = [=] __device__ { return 1; };

        nvstd::function<int()> fn1 = device_function; // ERROR, initialized with address of __device__ function
        nvstd::function<int()> fn2 = lambda_device;   // ERROR, initialized with address of functor with
                                                      //        __device__ operator() function
        kernel<<<1, 1>>>(fn2);                        // ERROR, passing nvstd::function from host to device
    }

    namespace nvstd {

    template <typename RetType, typename ...ArgTypes>
    class function<RetType(ArgTypes...)> {
    public:
        // constructors
        __device__ __host__ function() noexcept;
        __device__ __host__ function(nullptr_t) noexcept;
        __device__ __host__ function(const function&);
        __device__ __host__ function(function&&);

        template<typename F>
        __device__ __host__ function(F);

        // destructor
        __device__ __host__ ~function();

        // assignment operators
        __device__ __host__ function& operator=(const function&);
        __device__ __host__ function& operator=(function&&);
        __device__ __host__ function& operator=(nullptr_t);
        template<typename F>
        __device__ __host__ function& operator=(F&&);

        // swap
        __device__ __host__ void swap(function&) noexcept;

        // function capacity
        __device__ __host__ explicit operator bool() const noexcept;

        // function invocation
        __device__ RetType operator()(ArgTypes...) const;
    };

    // null pointer comparisons
    template <typename R, typename... ArgTypes>
    __device__ __host__
    bool operator==(const function<R(ArgTypes...)>&, nullptr_t) noexcept;

    template <typename R, typename... ArgTypes>
    __device__ __host__
    bool operator==(nullptr_t, const function<R(ArgTypes...)>&) noexcept;

    template <typename R, typename... ArgTypes>
    __device__ __host__
    bool operator!=(const function<R(ArgTypes...)>&, nullptr_t) noexcept;

    template <typename R, typename... ArgTypes>
    __device__ __host__
    bool operator!=(nullptr_t, const function<R(ArgTypes...)>&) noexcept;

    // specialized algorithms
    template <typename R, typename... ArgTypes>
    __device__ __host__
    void swap(function<R(ArgTypes...)>&, function<R(ArgTypes...)>&);

    } // namespace nvstd

    namespace cuda {   // same for "nv" and "cooperative_groups" namespaces

    struct foo;        // ERROR, class declaration in the "cuda" namespace

    void bar();        // ERROR, function declaration in the "cuda" namespace

    namespace utils {} // ERROR, namespace declaration in the "cuda" namespace

    } // namespace cuda

    namespace utils {
    namespace cuda {

    // CORRECT, namespace "cuda" may be used nested within a non-reserved namespace
    void bar();

    } // namespace cuda
    } // namespace utils

    // ERROR, Equivalent to adding symbols to namespace "cuda" at global scope
    using namespace utils;

    __host__ void host_function() {
        int x;                   // CORRECT, __host__ variable
        __device__   int y;      // ERROR,   __device__ variable declaration within a host function
        __tile__     int z;      // ERROR,   __tile__ variable declaration within a host function
        __shared__   int w;      // ERROR,   __shared__ variable declaration within a host function
        __managed__  int h;      // ERROR,   __managed__ variable  declaration within a host function
        __constant__ int i;      // ERROR,   __constant__ variable declaration within a host function
        extern __device__ int j; // CORRECT, extern __device__ variable
    }

    __device__ void device_function() {
        int x;                   // CORRECT, __device__ variable
        __constant__      int y; // ERROR,   __constant__ variable declaration within a device function
        __managed__       int z; // ERROR,   __managed__ variable  declaration within a device function
        extern __device__ int k; // CORRECT, extern __device__ variable
    }

    const            int   ConstVar          = 10;
    const            float ConstFloatVar     = 5.0f;
    inline constexpr float ConstexprFloatVar = 5.0f; // C++17

    struct MyStruct {
        static const            int   ConstVar          = 20;
    //  static const             float ConstFloatVar     = 5.0f; // ERROR, static const variables cannot be float
        static inline constexpr float ConstexprFloatVar = 5.0f; // CORRECT
    };

    extern const int ExternVar;

    __device__ void foo() {
        int array1[ConstVar];                     // CORRECT
        int array2[MyStruct::ConstVar];           // CORRECT

        const     float var1 = ConstFloatVar;     // CORRECT, except when the host compiler is Microsoft Visual Studio.
        constexpr float var2 = ConstexprFloatVar; // CORRECT
    //  int             var3 = ExternVar;          // ERROR, "ExternVar" is not initialized with a constant expression
    //  int&            var4 = ConstVar;           // ERROR, reference to host variable
    //  int*            var5 = &ConstVar;          // ERROR, address of host variable
    }

    __global__ void kernel(int* mmio_reg0, int* mmio_reg1) {
        // Write to MMIO register:
        int value = 13;
        asm volatile("st.relaxed.mmio.sys.u32 [%0], %1;"
            :
            : "l"(mmio_reg0), "r"(value) : "memory");

        // Read MMIO register:
        asm volatile("ld.relaxed.mmio.sys.u32 %0, [%1];"
            : "=r"(value)
            : "l"(mmio_reg1) : "memory");

        if (value != 42)
            __trap(); // Errors if wrong data read
    }

    struct TrivialStruct {
        int x;
    };

    struct NonTrivialStruct {
        __device__ NonTrivialStruct(int x) {}
    };

    __device__ void device_function(int x) {
        static int v1;              // CORRECT, implicit __device__ memory space specifier
        static int v2 = 11;         // CORRECT, implicit __device__ memory space specifier
    //  static int v3 = x;           // ERROR, dynamic initialization is not allowed

        static __managed__  int v4; // CORRECT, explicit
        static __device__   int v5; // CORRECT, explicit
        static __constant__ int v6; // CORRECT, explicit
        static __shared__   int v7; // CORRECT, explicit

        static TrivialStruct    s1;     // CORRECT, implicit __device__ memory space specifier
        static TrivialStruct    s2{22}; // CORRECT, implicit __device__ memory space specifier
    //  static TrivialStruct    s3{x};   // ERROR, dynamic initialization is not allowed
    //  static NonTrivialStruct s4{3};   // ERROR, dynamic initialization is not allowed
    }

    __host__ __device__ void host_device_function() {
        static            int v1; // CORRECT, implicit __device__ memory space specifier
    //  static __device__ int v2;  // ERROR, __device__-only variable inside a host-device function
    #ifdef __CUDA_ARCH__
        static __device__ int v3; // CORRECT, declaration is only visible during device compilation
    #else
        static int v4;            // CORRECT, declaration is only visible during host compilation
    #endif
    }

    #include <cassert>

    __host__ __device__ int host_device_function() {
        static int v = 0;
        v++;
        return v;
    }

    __global__ void kernel() {
        int ret = host_device_function(); // v = 1
        assert(ret == 4);                 // FAIL
    }

    int main() {
        host_device_function();           // v = 1
        host_device_function();           // v = 2
        int ret = host_device_function(); // v = 3
        assert(ret == 3);                 // OK
        kernel<<<1, 1>>>();
        cudaDeviceSynchronize();
    }

    __device__        int x; // OK
    extern __device__ int y; // ERROR in whole program compilation mode
    extern __shared__ int z; // OK

    //first.cu:
    struct S;                   // forward declaration
    __device__ void foo(S);     // ERROR, type 'S' is an incomplete type
    __device__ auto* ptr = foo; // ODR-use, address taken

    int main() {}

    //second.cu:
    struct S {};               // struct definition
    __device__ void foo(S) {}  // function definition

    void device_function1(__device__ int x) { } // ERROR, __device__ parameter
    void device_function2(__shared__ int x) { } // ERROR, __shared__ parameter

    #include <cassert>

    struct MyStruct {
        int  value = 1;
        int* ptr;

        MyStruct() = default;

        __host__ __device__ MyStruct(const MyStruct&) { ptr = &value; }
    };

    __global__ void device_function(MyStruct my_struct) {
        // this assert fails because "my_struct" is obtained by copying
        // the raw memory content and the copy constructor is skipped.
        assert(my_struct.ptr == &my_struct.value); // FAIL
    }

    void host_function(MyStruct my_struct) {
        assert(my_struct.ptr == &my_struct.value); // CORRECT
    }

    int main() {
        MyStruct my_struct;
        host_function(my_struct);
        device_function<<<1, 1>>>(my_struct); // copy constructor invoked in the host-side only
        cudaDeviceSynchronize();
    }

    #include <cassert>

    __managed__ int var = 0;

    struct MyStruct {
        __host__ __device__ ~MyStruct() { var = 3; }
    };

    __global__ void device_function(MyStruct my_struct) {
        assert(var == 0); // FAIL, MyStruct::~MyStruct() sets the value to 3
    }

    int main() {
        MyStruct my_struct;
        // GPU kernel execution is asynchronous with host execution.
        // As a result, MyStruct::~MyStruct() could be executed before
        // the kernel finishes executing.
        device_function<<<1, 1>>>(my_struct);
        cudaDeviceSynchronize();
    }

    struct MyStruct {
       static inline constexpr int value1 = 10; // C++17
       static constexpr        int value2 = 10; // C++11
       static const            int value3 = 10;
    // static                  int value4; // ERROR
    };

    struct MyStruct {
        friend __global__ void f();   // CORRECT, friend declaration only

    //  friend __global__ void g() {} // ERROR, friend definition
    };

    class Base {
        int x;
    public:
        __host__ __device__ Base() : x(10) {}
    };

    class Derived : public Base {
        int y;
    };

    class Other: public Base {
        int z;
    };

    __device__ void foo() {
        Derived D1;
        Other D2;
    }

    __host__ void bar() {
        Other D3;
    }

    struct Base1 {
        virtual __host__ __device__ ~Base1() {}
    };

    struct Derived1 : Base1 {}; // implicitly-declared virtual destructor
                                // ~Derived1() has __host__ __device__  execution space specifiers

    struct Base2 {
        virtual __device__ ~Base2() = default;
    };

    struct Derived2 : Base2 {}; // implicitly-declared virtual destructor
                                // ~Derived2() has __device__ execution space specifiers

    struct MyClass {
        virtual __host__ __device__ void f() {}
    };

    __global__ void kernel(MyClass my_class) {
        my_class.f(); // undefined behavior
    }

    int main() {
        MyClass my_class;
        kernel<<<1, 1>>>(my_class);
        cudaDeviceSynchronize();
    }

    struct BaseClass {
        virtual __host__ __device__ void f() {}
    };

    struct DerivedClass : BaseClass {
        __device__ void f() override {} // ERROR
    };

    template <typename T>
    __global__ void kernel() {}

    template <typename T>
    __device__ int device_var; // C++14

    struct {
        int v;
    } unnamed_struct;

    void host_function() {
        struct LocalStruct {};
    //  kernel<LocalStruct><<<1, 1>>>(); // ERROR, LocalStruct is defined within a host function
        int data = 4;
    //  cudaMemcpyToSymbol(device_var<LocalStruct>, &data, sizeof(data)); // ERROR, same as above

        auto lambda = [](){};
    //  kernel<decltype(lambda)><<<1, 1>>>();         // ERROR, unnamed type
    //  kernel<decltype(unnamed_struct)><<<1, 1>>>(); // ERROR, unnamed type
    }

    class MyClass {
    private:
        struct PrivateStruct {};
    public:
        static void launch() {
    //      kernel<PrivateStruct><<<1, 1>>>(); // ERROR, private type
        }
    };

    __device__ int my_var; // global scope

    inline namespace NS {

    __device__ int my_var; // namespace scope

    } // namespace NS

    constexpr int ConstexprVar = 4; // scalar type

    struct MyStruct {
        static constexpr int ConstexprVar = 100;
    };

    constexpr MyStruct my_struct = MyStruct{}; // class type

    constexpr int array[] = {1, 2, 3};

    __device__ constexpr int get_value(int idx) {
        return array[idx];                      // CORRECT
    }

    __device__ void foo(int idx) {
        int        v1 = ConstexprVar;           // CORRECT
        int        v2 = MyStruct::ConstexprVar; // CORRECT
    //  const int &v3 = ConstexprVar1;          // ERROR, reference to host constexpr variable
    //  const int *v4 = &ConstexprVar1;         // ERROR, address of host constexpr variable
        int        v5 = get_value(2);           // CORRECT, 'get_value(2)' is a constant expression.
    //  int        v6 = get_value(idx);         // ERROR, 'get_value(idx)' is not a constant expression
    //  int        v7 = array[2];               // ERROR, 'array' is not scalar type.
        MyStruct   v8 = my_struct;              // CORRECT
    }

    template <typename... Pack>
    __global__ void kernel1(); // CORRECT

    // template <typename... Pack, template T>
    // __global__ void kernel2(); // ERROR, parameter pack is not the last parameter

    template <typename... TArgs>
    struct MyStruct {};

    // template <typename... Pack1, typename... Pack2>
    // __global__ void kernel3(MyStruct<Pack1...>, MyStruct<Pack2...>); // ERROR, more than one parameter pack

    struct MyStruct1 {
        MyStruct1() = default;
    };

    void host_function() {
        MyStruct1 my_struct; // __host__ __device__ constructor
    }

    __device__ void device_function() {
        MyStruct1 my_struct; // __host__ __device__ constructor
    }

    struct MyStruct2 {
        __device__ MyStruct2() = default; // WARNING: __device__ annotation is ignored
    };

    struct MyStruct3 {
        __host__ MyStruct3();
    };
    MyStruct3::MyStruct3() = default; // out-of-line definition, not ignored

    __device__ void device_function2() {
    //  MyStruct3 my_struct; // ERROR, __host__ constructor
    }

    struct MyStruct4 {
        //  MyStruct4::~MyStruct4 has host execution space, not ignored because virtual
        virtual __host__ ~MyStruct4() = default;
    };

    __device__ void device_function3() {
        MyStruct4 my_struct4;
        // implicit destructor call for 'my_struct4':
        //    ERROR: call from a __device__ function 'device_function3' to a
        //    __host__ function 'MyStruct4::~MyStruct4'
    }

    #include <initializer_list>

    __device__ void foo(std::initializer_list<int> in) {}

    __device__ void bar() {
        foo({4,5,6}); // (a) initializer list containing only constant expressions.
        int i = 4;
        foo({i,5,6}); // (b) initializer list with at least one  non-constant element.
                      // This form may have better performance than (a).
    }

     __device__ auto device_function(int x) { // deduced return type
         return x;                            // decltype(auto) has the same behavior
     }

     __global__ void kernel() {
         int x = sizeof(device_function(2));         // CORRECT, device code scope
     }

     // const int size = sizeof(device_function(2)); // ERROR, return type deduction on host

     void host_function() {
     //  using T = decltype(device_function(2));     // ERROR, return type deduction on host
     }

    void host_fn1() {
      // ERROR, referenced outside device function bodies
      int (*p1)(int) = fn1;

      struct S_local_t {
        // ERROR, referenced outside device function bodies
        decltype(fn2(10)) m1;

        S_local_t() : m1(10) { }
      };
    }

    // ERROR, referenced outside device function bodies
    template <typename T = decltype(fn2)>
    void host_fn2() { }

    template<typename T> struct MyStruct { };

    // ERROR, referenced outside device function bodies
    struct S1_derived_t : MyStruct<decltype(fn1)> { };

    // ERROR on Windows (non-portable), const-qualified
    template <typename T>
    __device__ const T var = 0;

     // CORRECT, ptr1 is not const-qualified
    template <typename T>
    __device__ const T* ptr1 = nullptr;

    // ERROR on Windows (non-portable), ptr2 is const-qualified
    template <typename T>
    __device__ const T* const ptr2 = nullptr;

    inline        __device__ int device_var1;  // CORRECT, when compiled in Separate Compilation mode (-rdc=true or -dc)
                                               // ERROR, when compiled in Whole Program Compilation mode

    static inline __device__ int device_var2;  // CORRECT, internal linkage

    namespace {

    inline __device__ int device_var3;         // CORRECT, internal linkage

    inline __shared__ int shared_var;          // CORRECT, internal linkage

    static inline __device__ int device_var4;  // CORRECT, internal linkage

    inline __device__ int device_var5;         // CORRECT, internal linkage

    } // namespace

    struct S {
        int x, y;
    };
    // __device__ auto [a, b] = S{4, 5}; // ERROR

    #include <compare> // std::strong_ordering implementation

    struct S {
        int x, y;

        auto operator<=>(const S&) const = default; // (a)

        __host__ __device__ bool operator<=>(int rhs) const { return false; } // (b)
    };

    __host__ __device__ bool host_device_function(S a, S b) {
        if (a <=> 1)  // CORRECT, calls a user-defined host-device overload (b)
            return true;
        return a < b; // CORRECT, call to an implicitly-declared function (a)
                      // Note: it requires a device-compatible std::strong_ordering
                      //       implementation provided in the header <compare>
                      //       and the flag --expt-relaxed-constexpr
    }

    consteval int host_consteval() {
        return 10;
    }

    __device__ consteval int device_consteval() {
        return 10;
    }

    __device__ int device_function() {
        return host_consteval();   // CORRECT, even if called from device code
    }

    __host__ __device__ int host_device_function() {
        return device_function();  // CORRECT, even if called from host-device code
    }