5.4. C/C++ 语言扩展（C/C++ Language Extensions）

    __host__ __device__ void func() {
    #if defined(__CUDA_ARCH__)
        // Device code path
    #else
        // Host code path
    #endif
    }

    __device__   float device_var       = 4.0f; // Variable in device memory
    __constant__ float constant_mem_var = 4.0f; // Variable in constant memory
                                                // For readability, the following example focuses on a device variable.
    int main() {
        float* device_ptr;
        cudaGetSymbolAddress((void**) &device_ptr, device_var);        // Gets address of device_var

        size_t symbol_size;
        cudaGetSymbolSize(&symbol_size, device_var);                   // Retrieves the size of the symbol (4 bytes).

        float host_var;
        cudaMemcpyFromSymbol(&host_var, device_var, sizeof(host_var)); // Copies from device to host.

        host_var = 3.0f;
        cudaMemcpyToSymbol(device_var, &host_var, sizeof(host_var));   // Copies from host to device.
    }

    extern __shared__ char dynamic_smem_pointer[];
    // extern __shared__ char* dynamic_smem_pointer; alternative syntax

    __global__ void kernel() { // or a __device__ function
        __shared__ int smem_var1[4];                  // static size
        auto smem_var2 = (int*) dynamic_smem_pointer; // dynamic size
    }

    int main() {
        size_t shared_memory_size = 16;
        kernel<<<1, 1, shared_memory_size>>>();
        cudaDeviceSynchronize();
    }

    #include <cassert>

    __device__ __managed__ int global_var = 10; // OK

    int* ptr = &global_var;                     // ERROR: use of a managed variable in static initialization

    struct MyStruct1 {
        int field;
        MyStruct1() : field(global_var) {};
    };

    struct MyStruct2 {
        ~MyStruct2() { global_var = 10; }
    };

    MyStruct1 temp1; // ERROR: use of managed variable in dynamic initialization

    MyStruct2 temp2; // ERROR: use of managed variable in the destructor of
                     //        object with static storage duration

    __device__ __managed__ const int const_var = 10;         // ERROR: const-qualified type

    __device__ __managed__ int&      reference = global_var; // ERROR: reference type

    template <int* Addr>
    struct MyStruct3 {};

    MyStruct3<&global_var> temp;     // ERROR: address of managed variable is not a constant expression

    __global__ void kernel(int* ptr) {
        assert(ptr == &global_var);  // OK
        global_var = 20;             // OK
    }

    int main() {
        int* ptr = &global_var;      // OK
        kernel<<<1, 1>>>(ptr);
        cudaDeviceSynchronize();
        global_var++;                // OK
        decltype(global_var) var1;   // ERROR: managed variable used as unparenthesized argument to decltype

        decltype((global_var)) var2; // OK
    }

    __device__ int x_device = 0;
    __tile__   int x_tile = 0;

    // ERROR: Variable cannot be both device and tile
    __device__ __tile__ int x_device_tile = 0;

    __global__ void simt_function(int* in) {
      x_device = 1; // OK: __device__ variable accessed from SIMT code
      x_tile = 1;   // ERROR: __tile__ variable accessed from SIMT code

      *in = 1; // OK: May point to any device global memory, including x_tile.
    }

    __tile_global__ void tile_function(int* in) {
      x_device = 1; // ERROR: __device__ variable accessed from tile code
      x_tile = 1;   // OK: __tile__ variable accessed from tile code

      *in = 1; // OK: May point to any device global memory, including x_device.
    }

    int main() {
      int* x_device_ptr;
      cudaGetSymbolAddress((void**) &x_device_ptr, x_device);    // Gets address of x_device

      int* x_tile_ptr;
      cudaGetSymbolAddress((void**) &x_tile_ptr, x_tile);        // Gets address of x_tile

      // OK: Passing pointer to tile variable into SIMT kernel
      simt_function<<<1,1>>>(x_tile_ptr);
      cudaDeviceSynchronize();

      // OK: Passing pointer to device variable into tile kernel
      tile_function<<<1,1>>>(x_device_ptr);
      cudaDeviceSynchronize();
    }

    __tile__ int* ptr; // ERROR

    struct S1 { int* ptr; };
    __tile__ S1 val; // ERROR

    __device__
    void device_function(const float* a, const float* b, float* c) {
        c[0] = a[0] * b[0];
        c[1] = a[0] * b[0];
        c[2] = a[0] * b[0] * a[1];
        c[3] = a[0] * a[1];
        c[4] = a[0] * b[0];
        c[5] = b[0];
        ...
    }

    __device__
    void device_function(const float* __restrict__ a, const float* __restrict__ b, float* __restrict__ c);

    __device__
    void device_function(const float* __restrict__ a, const float* __restrict__ b, float* __restrict__ c) {
        float t0 = a[0];
        float t1 = b[0];
        float t2 = t0 * t1;
        float t3 = a[1];
        c[0]     = t2;
        c[1]     = t2;
        c[4]     = t2;
        c[2]     = t2 * t3;
        c[3]     = t0 * t3;
        c[5]     = t1;
        ...
    }

    __global__
    void kernel1(const float* in, float* out) {
        *out = *in; // PTX: ld.global
    }

    __global__
    void kernel2(const float* __restrict__ in, float* out) {
        *out = *in;  // PTX: ld.global.nc
    }

    struct MyStruct {
        int         x;
        mutable int y;
    };

    __device__ void external_function(const MyStruct&);

    __global__ void kernel(const __grid_constant__ MyStruct s) {
        // s.x++; // Compile error: tried to modify read-only memory
        // s.y++; // Undefined Behavior: tried to modify read-only memory

        // Compiler will NOT create a per-thread local copy of "s":
        external_function(s);
    }

    int sum(int4 value) {
        return value.x + value.y + value.z + value.w;
    }

    int4 add_one(int x, int y, int z, int w) {
        return make_int4(x + 1, y + 1, z + 1, w + 1);
    }

    __global__ void kernel(float* parameter);

    kernel<<<grid_dim, block_dim, dynamic_smem_bytes>>>(parameter);

    __global__ void __cluster_dims__(2, 1, 1) kernel(float* parameter);

    __global__ void kernel(float parameter1, int parameter2) {}

    int main() {
        cudaLaunchConfig_t config = {0};
        // The grid dimension is not affected by cluster launch, and is still enumerated
        // using the number of blocks.
        // The grid dimension should be a multiple of cluster size.
        config.gridDim          = dim3{4};  // 4 blocks
        config.blockDim         = dim3{32}; // 32 threads per block
        config.dynamicSmemBytes = 1024;     // 1 KB

        cudaLaunchAttribute attribute[1];
        attribute[0].id               = cudaLaunchAttributeClusterDimension;
        attribute[0].val.clusterDim.x = 2; // Cluster size in X-dimension
        attribute[0].val.clusterDim.y = 1;
        attribute[0].val.clusterDim.z = 1;
        config.attrs    = attribute;
        config.numAttrs = 1;

        float parameter1 = 3.0f;
        int   parameter2 = 4;
        cudaLaunchKernelEx(&config, kernel, parameter1, parameter2);
    }

    __global__ void
    __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor, maxBlocksPerCluster)
    MyKernel(...) {
        ...
    }

    #define THREADS_PER_BLOCK  256

    #if __CUDA_ARCH__ >= 900
        #define MY_KERNEL_MAX_THREADS  (2 * THREADS_PER_BLOCK)
        #define MY_KERNEL_MIN_BLOCKS   3
    #else
        #define MY_KERNEL_MAX_THREADS  THREADS_PER_BLOCK
        #define MY_KERNEL_MIN_BLOCKS   2
    #endif

    __global__ void
    __launch_bounds__(MY_KERNEL_MAX_THREADS, MY_KERNEL_MIN_BLOCKS)
    MyKernel(...) {
        ...
    }

    // Host code
    MyKernel<<<blocksPerGrid, MY_KERNEL_MAX_THREADS>>>(...);

    // Host code
    MyKernel<<<blocksPerGrid, THREADS_PER_BLOCK>>>(...);

    // Host code
    cudaGetDeviceProperties(&deviceProp, device);
    int threadsPerBlock = (deviceProp.major >= 9) ? 2 * THREADS_PER_BLOCK : THREADS_PER_BLOCK;
    MyKernel<<<blocksPerGrid, threadsPerBlock>>>(...);

    __global__ void
    __maxnreg__(maxNumberRegistersPerThread)
    MyKernel(...) {
        ...
    }

    void __syncthreads();
    int  __syncthreads_count(int predicate);
    int  __syncthreads_and(int predicate);
    int  __syncthreads_or(int predicate);

    #include <cuda_runtime_api.h>
    #include <memory.h>
    #include <cstdlib>
    #include <ctime>
    #include <stdio.h>


    // assuming blockDim.x is 128
    __global__ void example_syncthreads(int* input_data, int* output_data)
    {
        __shared__ int shared_data[128];
        shared_data[threadIdx.x] = input_data[blockDim.x*blockIdx.x + threadIdx.x];

        // All threads synchronize, guaranteeing all writes to 'shared_data' are ordered
        // before any thread is unblocked from '__syncthreads()':
        __syncthreads();

        // A single thread safely reads 'shared_data':
        if (threadIdx.x == 0) {
            float sum = 0;
            for (int i = 0; i < blockDim.x; ++i) {
                sum += shared_data[i];
            }
            output_data[blockIdx.x] = sum;
        }
    }


    void initArray(int* A, int length)
    {
         std::srand(std::time({}));
        for(int i=0; i<length; i++)
        {
            A[i] = int(10.f * (rand() / (float)RAND_MAX));
        }
    }


    int main(int argc, char** arg)
    {
        constexpr int block_size = 128;
        constexpr int input_length = 1024;
        constexpr int output_length = input_length / block_size;

        // Pointers to memory vectors
        int* input = nullptr;
        int* output = nullptr;
        int* comparisonResult = (int*)malloc(output_length*sizeof(int));

        // Use unified memory to allocate buffers
        cudaMallocManaged(&input, input_length*sizeof(int));
        cudaMallocManaged(&output, output_length*sizeof(int));

        initArray(input, input_length);
        int grid_size = input_length / block_size;

        example_syncthreads<<<grid_size, block_size>>>(input, output);
        cudaDeviceSynchronize();

        for(int i=0; i<output_length; i++)
        {
            comparisonResult[i] = 0;
            for(int j=0; j < block_size; j++)
            {
                comparisonResult[i] += input[i*block_size + j];
            }
        }

        for(int i=0; i< output_length; i++)
        {
            if(output[i] != comparisonResult[i])
            {
                printf("Results do not match at index %d: %d != %d\n", i, output[i], comparisonResult[i]);
                exit(-1);
            }
        }
        printf("Test passed\n");

        return 0;
    }

    // assuming blockDim.x is 128
    __global__ void syncthreads_valid_behavior(int* input_data, int* output_data) {
        __shared__ int shared_data[128];
        shared_data[threadIdx.x] = input_data[threadIdx.x];
        if (blockIdx.x > 0) { // CORRECT, uniform condition across all block threads
            __syncthreads();
            output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
        }
    }

    // assuming blockDim.x is 128
    __global__ void syncthreads_invalid_behavior1(int* input_data, int* output_data) {
        __shared__ int shared_data[256];
        shared_data[threadIdx.x] = input_data[threadIdx.x];
        if (threadIdx.x > 0) { // WRONG, non-uniform condition
            __syncthreads();   // Undefined Behavior
            output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
        }
    }

    // assuming blockDim.x is 128
    __global__ void syncthreads_invalid_behavior2(int* input_data, int* output_data) {
        __shared__ int shared_data[256];
        shared_data[threadIdx.x] = input_data[threadIdx.x];
        for (int i = 0; i < blockDim.x; ++i) {
            if (i == threadIdx.x) { // WRONG, non-uniform condition
                __syncthreads();    // Undefined Behavior
            }
        }
        output_data[threadIdx.x] = shared_data[128 - threadIdx.x];
    }

    int __syncthreads_count(int predicate);

    int __syncthreads_and(int predicate);

    int __syncthreads_or(int predicate);

    void __syncwarp(unsigned mask = 0xFFFFFFFF);

    __global__ void example_syncwarp(int* input_data, int* output_data) {
        if (threadIdx.x < warpSize) {
            __shared__ int shared_data[warpSize];
            shared_data[threadIdx.x] = input_data[threadIdx.x];

            __syncwarp(); // equivalent to __syncwarp(0xFFFFFFFF)
            if (threadIdx.x == 0)
                output_data[0] = shared_data[1];
        }
    }

    __device__ int X = 1, Y = 2;

    __device__ void writeXY() {
        X = 10;
        Y = 20;
    }

    __device__ void readXY() {
        int B = Y;
        int A = X;
    }

    // <cuda/atomic> header
    cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_block);

    void __threadfence_block();

    cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);

    void __threadfence();

    cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_system);

    void __threadfence_system();

    #include <cuda/atomic>

    __device__ int X = 1, Y = 2;

    __device__ void writeXY() {
        X = 10;
        cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);
        Y = 20;
    }

    __device__ void readXY() {
        int B = Y;
        cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);
        int A = X;
    }

    __device__ int X = 1, Y = 2;

    __device__ void writeXY() {
        X = 10;
        __threadfence();
        Y = 20;
    }

    __device__ void readXY() {
        int B = Y;
        __threadfence();
        int A = X;
    }

    #include <cuda/atomic>

    __device__ int count = 0;

    __global__ void sum(const float*    array,
                        int             N,
                        volatile float* result) {
        __shared__ bool isLastBlockDone;
        // Each block sums a subset of the input array.
        float partialSum = calculatePartialSum(array, N);

        if (threadIdx.x == 0) {
            // Thread 0 of each block stores the partial sum to global memory.
            // The compiler will use a store operation that bypasses the L1 cache
            // since the "result" variable is declared as volatile.
            // This ensures that the threads of the last block will read the correct
            // partial sums computed by all other blocks.
            result[blockIdx.x] = partialSum;

            // Thread 0 makes sure that the increment of the "count" variable is
            // only performed after the partial sum has been written to global memory.
            cuda::atomic_thread_fence(cuda::memory_order_seq_cst, cuda::thread_scope_device);

            // Thread 0 signals that it is done.
            int count_old = atomicInc(&count, gridDim.x);

            // Thread 0 determines if its block is the last block to be done.
            isLastBlockDone = (count_old == (gridDim.x - 1));
        }
        // Synchronize to make sure that each thread reads the correct value of
        // isLastBlockDone.
        __syncthreads();

        if (isLastBlockDone) {
            // The last block sums the partial sums stored in result[0 .. gridDim.x-1]
            float totalSum = calculateTotalSum(result);

            if (threadIdx.x == 0) {
                // Thread 0 of last block stores the total sum to global memory and
                // resets the count variable, so that the next kernel call works
                // properly.
                result[0] = totalSum;
                count     = 0;
            }
        }
    }

    #include <cuda_runtime.h>

    __global__ void atomicAdd_kernel(int* addr) {
        atomicAdd_system(addr, 10);
    }

    void test_atomicAdd(int device_id) {
        int* addr;
        cudaMallocManaged(&addr, 4);
        *addr = 0;

        cudaDeviceProp deviceProp;
        cudaGetDeviceProperties(&deviceProp, device_id);
        if (deviceProp.concurrentManagedAccess != 1) {
            return; // the device does not coherently access managed memory concurrently with the CPU
        }

        atomicAdd_kernel<<<...>>>(addr);
        __sync_fetch_and_add(addr, 10);  // CPU atomic operation
    }

    #include <cuda/memory>
    #include <cuda/std/bit>

    __device__ float customAtomicAdd(float* d_ptr, float value) {
        volatile unsigned* d_ptr_unsigned = reinterpret_cast<unsigned*>(d_ptr);
        unsigned  old_value      = *d_ptr_unsigned;
        unsigned  assumed;
        do {
            assumed                          = old_value;
            float    assumed_float           = cuda::std::bit_cast<float>(assumed);
            float    expected_value          = assumed_float + value;
            unsigned expected_value_unsigned = cuda::std::bit_cast<unsigned>(expected_value);
            old_value                        = atomicCAS(d_ptr_unsigned, assumed, expected_value_unsigned);
        // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
        } while (assumed != old_value);
        return cuda::std::bit_cast<float>(old_value);
    }

    T atomicAdd(T* address, T val);

    T atomicSub(T* address, T val);

    unsigned atomicInc(unsigned* address, unsigned val);

    unsigned atomicDec(unsigned* address, unsigned val);

    T atomicAnd(T* address, T val);

    T atomicOr(T* address, T val);

    T atomicXor(T* address, T val);

    T atomicMin(T* address, T val);

    T atomicMax(T* address, T val);

    T atomicExch(T* address, T val);

    template<typename T>
    T atomicExch(T* address, T val); // only 128-bit types, compute capability 9.x and higher

    T atomicCAS(T* address, T compare, T val);

    template<typename T>
    T atomicCAS(T* address, T compare, T val);  // only 128-bit types, compute capability 9.x and higher

    // atomic memory orders
    enum {
       __NV_ATOMIC_RELAXED,
       __NV_ATOMIC_CONSUME,
       __NV_ATOMIC_ACQUIRE,
       __NV_ATOMIC_RELEASE,
       __NV_ATOMIC_ACQ_REL,
       __NV_ATOMIC_SEQ_CST
    };

    // thread scopes
    enum {
       __NV_THREAD_SCOPE_THREAD,
       __NV_THREAD_SCOPE_BLOCK,
       __NV_THREAD_SCOPE_CLUSTER,
       __NV_THREAD_SCOPE_DEVICE,
       __NV_THREAD_SCOPE_SYSTEM
    };

    __device__ T __nv_atomic_load_n(T*  pointer,
                                    int memory_order,
                                    int thread_scope = __NV_THREAD_SCOPE_SYSTEM);

     // Not permitted in a host function
     __host__ void bar() {
         unsigned u1 = 1, u2 = 2;
         __nv_atomic_load(&u1, &u2, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_SYSTEM);
     }

     // Not permitted to be applied to local memory
    __device__ void foo() {
       unsigned a = 1, b;
       __nv_atomic_load(&a, &b, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_SYSTEM);
    }

     // Not permitted as a template default argument.
     // The function address cannot be taken.
     template<void *F = __nv_atomic_load_n>
     class X {
         void *f = F; // The function address cannot be taken.
     };

     // Not permitted to be called in a constructor initialization list.
     class Y {
         int a;
     public:
         __device__ Y(int *b): a(__nv_atomic_load_n(b, __NV_ATOMIC_RELAXED)) {}
     };

    __device__ T    __nv_atomic_fetch_add(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_add      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_fetch_sub(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_sub      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_fetch_and(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_and      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_fetch_or(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_or      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_fetch_xor(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_xor      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_fetch_min(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_min      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_fetch_max(T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_max      (T* address, T val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ T    __nv_atomic_exchange_n(T* address, T val,          int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_exchange  (T* address, T* val, T* ret, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ bool __nv_atomic_compare_exchange  (T* address, T* expected, T* desired, bool weak, int success_order, int failure_order,
                                                   int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ bool __nv_atomic_compare_exchange_n(T* address, T* expected, T desired, bool weak, int success_order, int failure_order,
                                                   int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ void __nv_atomic_load  (T* address, T* ret, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ T    __nv_atomic_load_n(T* address,         int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ void __nv_atomic_store  (T* address, T* val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);
    __device__ void __nv_atomic_store_n(T* address, T  val, int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    __device__ void __nv_atomic_thread_fence(int order, int scope = __NV_THREAD_SCOPE_SYSTEM);

    unsigned __activemask();

    // Check whether at least one thread's predicate evaluates to true
    if (pred) {
        // Invalid: the value of 'at_least_one' is non-deterministic
        // and could vary between executions.
        at_least_one = __activemask() > 0;
    }

    unsigned mask      = __activemask();              // Assume mask == 0xFFFFFFFF (all bits set, all threads active)
    int      predicate = threadIdx.x % 2 == 0;        // 1 for even threads, 0 for odd threads
    int      result    = __any_sync(mask, predicate); // Active threads might not be preserved

    int      __all_sync   (unsigned mask, int predicate);
    int      __any_sync   (unsigned mask, int predicate);
    unsigned __ballot_sync(unsigned mask, int predicate);

    unsigned __match_any_sync(unsigned mask, T value);
    unsigned __match_all_sync(unsigned mask, T value, int *pred);

    T        __reduce_add_sync(unsigned mask, T value);
    T        __reduce_min_sync(unsigned mask, T value);
    T        __reduce_max_sync(unsigned mask, T value);

    unsigned __reduce_and_sync(unsigned mask, unsigned value);
    unsigned __reduce_or_sync (unsigned mask, unsigned value);
    unsigned __reduce_xor_sync(unsigned mask, unsigned value);

    T __shfl_sync     (unsigned mask, T value, int      srcLane,  int width=warpSize);
    T __shfl_up_sync  (unsigned mask, T value, unsigned delta,    int width=warpSize);
    T __shfl_down_sync(unsigned mask, T value, unsigned delta,    int width=warpSize);
    T __shfl_xor_sync (unsigned mask, T value, int      laneMask, int width=warpSize);

    int laneId = threadIdx.x % warpSize;
    int data   = ...

    // all warp threads get 'data' from lane 0
    int result1 = __shfl_sync(0xFFFFFFFF, data, 0);

    if (laneId < 4) {
        // lanes 0, 1, 2, 3 get 'data' from lane 1
        int result2 = __shfl_sync(0xb1111, data, 1);
    }

    // lanes [0 - 15] get 'data' from lane 0
    // lanes [16 - 31] get 'data' from lane 16
    int result3 = __shfl_sync(0xFFFFFFFF, value, warpSize / 2);

    // each lane gets 'data' from the lane two positions above
    // lanes 30, 31 get their original value
    int result4 = __shfl_down_sync(0xFFFFFFFF, data, 2);

    int laneId = threadIdx.x % warpSize;
    int value  = ...
     // undefined behavior: lane 0 does not participate in the call
    int result = (laneId > 0) ? __shfl_sync(0xFFFFFFFF, value, 0) : 0;

    if (laneId <= 4) {
        // undefined behavior: destination lanes 5, 6 are not active for lanes 3, 4
        result = __shfl_down_sync(0b11111, value, 2);
    }

    // undefined behavior: width is not a power of 2
    __shfl_sync(0xFFFFFFFF, value, 0, /*width=*/31);

    #include <cassert>
    #include <cuda/warp>

    __global__ void warp_broadcast_kernel(int input) {
        int laneId = threadIdx.x % 32;
        int value;
        if (laneId == 0) { // unused variable for all threads except lane 0
            value = input;
        }
        value = cuda::device::warp_shuffle_idx(value, 0); // Synchronize all threads in warp, and get "value" from lane 0
        assert(value == input);
    }

    int main() {
        warp_broadcast_kernel<<<1, 32>>>(1234);
        cudaDeviceSynchronize();
        return 0;
    }

    #include <assert.h>

    __global__ void warp_broadcast_kernel(int input) {
        int laneId = threadIdx.x % 32;
        int value;
        if (laneId == 0) { // unused variable for all threads except lane 0
            value = input;
        }
        value = __shfl_sync(0xFFFFFFFF, value, 0); // Synchronize all threads in warp, and get "value" from lane 0
        assert(value == input);
    }

    int main() {
        warp_broadcast_kernel<<<1, 32>>>(1234);
        cudaDeviceSynchronize();
        return 0;
    }

    #include <cstdio>
    #include <cub/cub.cuh>

    __global__ void scan_sub_partition_with_8_threads_kernel() {
        using WarpScan    = cub::WarpScan<int, 8>;
        using TempStorage = typename WarpScan::TempStorage;
        __shared__ TempStorage temp_storage;

        int laneId = threadIdx.x % 32;
        int value  = 31 - laneId; // starting value to accumulate
        int partial_sum;
        WarpScan(temp_storage).InclusiveSum(value, partial_sum);
        printf("Thread %d final value = %d\n", threadIdx.x, partial_sum);
    }

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

    #include <stdio.h>

    __global__ void scan_sub_partition_with_8_threads_kernel() {
        int laneId = threadIdx.x % 32;
        int value  = 31 - laneId; // starting value to accumulate
        // Loop to accumulate scan within my partition.
        // Scan requires log2(8) == 3 steps for 8 threads
        for (int delta = 1; delta <= 4; delta *= 2) {
            int tmp         = __shfl_up_sync(0xFFFFFFFF, value, delta, /*width=*/8); // read from laneId - delta
            int source_lane = laneId % 8 - delta;
            if (source_lane >= 0) // lanes with 'source_lane < 0' have their value unchanged
                value += tmp;
        }
        printf("Thread %d final value = %d\n", threadIdx.x, value);
    }

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

    #include <cstdio>
    #include <cub/cub.cuh>
    #include <cuda/warp>

    __global__ void warp_reduce_kernel() {
        using WarpReduce  = cub::WarpReduce<int>;
        using TempStorage = typename WarpReduce::TempStorage;
        __shared__ TempStorage temp_storage;

        int laneId     = threadIdx.x % 32;
        int value      = 31 - laneId; // starting value to accumulate
        auto aggregate = WarpReduce(temp_storage).Sum(value);
        aggregate      = cuda::device::warp_shuffle_idx(aggregate, 0);
        printf("Thread %d final value = %d\n", threadIdx.x, aggregate);
    }

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

    #include <stdio.h>

    __global__ void warp_reduce_kernel() {
        int laneId = threadIdx.x % 32;
        int value  = 31 - laneId; // starting value to accumulate
        // Use XOR mode to perform butterfly reduction
        // A full-warp reduction requires log2(32) == 5 steps
        for (int i = 1; i <= 16; i *= 2)
            value += __shfl_xor_sync(0xFFFFFFFF, value, i);
        // "value" now contains the sum across all threads
        printf("Thread %d final value = %d\n", threadIdx.x, value);
    }

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

    __global__ void valid_examples() {
        if (threadIdx.x < 4) {        // threads 0, 1, 2, 3 are active
            __all_sync(0b1111, pred); // CORRECT, threads 0, 1, 2, 3 participate in the call
        }

        if (threadIdx.x == 0)
            return; // exit
        // CORRECT, all non-exited threads participate in the call
        __all_sync(0xFFFFFFFF, pred);
    }

    __global__ void example_syncwarp_with_mask(int* input_data, int* output_data) {
        if (threadIdx.x < warpSize) {
            __shared__ int shared_data[warpSize];
            shared_data[threadIdx.x] = input_data[threadIdx.x];

            unsigned mask = threadIdx.x < 16 ? 0xFFFF : 0xFFFF0000; // CORRECT
            __syncwarp(mask);
            if (threadIdx.x == 0 || threadIdx.x == 16)
                output_data[threadIdx.x] = shared_data[threadIdx.x + 1];
        }
    }

    __global__ void example_syncwarp_with_mask_branches(int* input_data, int* output_data) {
        if (threadIdx.x < warpSize) {
            __shared__ int shared_data[warpSize];
            shared_data[threadIdx.x] = input_data[threadIdx.x];

            if (threadIdx.x < 16) {
                unsigned mask = 0xFFFF; // CORRECT
                __syncwarp(mask);
                output_data[threadIdx.x] = shared_data[15 - threadIdx.x];
            }
            else {
                unsigned mask = 0xFFFF0000; // CORRECT
                __syncwarp(mask);
                output_data[threadIdx.x] = shared_data[31 - threadIdx.x];
            }
        }
    }

    if (threadIdx.x < 4) {           // threads 0, 1, 2, 3 are active
        __all_sync(0b0000011, pred); // WRONG, threads 2, 3 are active but not set in mask
        __all_sync(0b1111111, pred); // WRONG, threads 4, 5, 6 are not active but set in mask
    }

    // WRONG, participating threads have a different and overlapping mask
    __all_sync(threadIdx.x == 0 ? 1 : 0xFFFFFFFF, pred);

    #if !defined(__CUDA_ARCH__)
        typedef int my_type;
    #else
        typedef double my_type;
    #endif

    __device__ my_type my_var;           // ERROR: my_var's type depends on __CUDA_ARCH__

    __global__ void kernel(my_type in) { // ERROR: kernel's type depends on __CUDA_ARCH__
        ...
    }

    __device__ int result;

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

    __host__ __device__ void host_device_function(void) {
    #if !defined(__CUDA_ARCH__)
        kernel<<<1, 1>>>(1); // ERROR: "kernel<int>" instantiation only
                                //        when __CUDA_ARCH__ is undefined!
    #endif
    }

    int main(void) {
        host_device_function();
        cudaDeviceSynchronize();
        return 0;
    }

    #if !defined(__CUDA_ARCH__)
        void host_function(void) {} // ERROR: The definition of host_function()
                                    //        is only present when __CUDA_ARCH__
                                    //        is undefined
    #endif

    template<typename T>
    __device__ T* get_ptr() {
    #if __CUDA_ARCH__ == 900
        return nullptr; /* no address */
    #else
        __shared__ T arr[256];
        return arr;
    #endif
    }

    __device__ __nv_pure__
    int add(int a, int b) {
        return a + b;
    }

    __device__ unsigned __isGlobal      (const void* ptr);
    __device__ unsigned __isShared      (const void* ptr);
    __device__ unsigned __isConstant    (const void* ptr);
    __device__ unsigned __isGridConstant(const void* ptr);
    __device__ unsigned __isLocal       (const void* ptr);

    __device__ size_t __cvta_generic_to_global  (const void* ptr); // PTX: cvta.to.global
    __device__ size_t __cvta_generic_to_shared  (const void* ptr); // PTX: cvta.to.shared
    __device__ size_t __cvta_generic_to_constant(const void* ptr); // PTX: cvta.to.const
    __device__ size_t __cvta_generic_to_local   (const void* ptr); // PTX: cvta.to.local

    __device__ void* __cvta_global_to_generic  (size_t raw_ptr); // PTX: cvta.global
    __device__ void* __cvta_shared_to_generic  (size_t raw_ptr); // PTX: cvta.shared
    __device__ void* __cvta_constant_to_generic(size_t raw_ptr); // PTX: cvta.const
    __device__ void* __cvta_local_to_generic   (size_t raw_ptr); // PTX: cvta.local

    __shared__ int smem_var;
    smem_var        = 42;
    size_t smem_ptr = __cvta_generic_to_shared(&smem_var);
    int    output;
    asm volatile("ld.shared.s32 %0, [%1];" : "=r"(output) : "l"(smem_ptr) : "memory");
    assert(output == 42);

    __shared__ int smem_var;
    uint32_t       smem_ptr_32bit = static_cast<uint32_t>(__cvta_generic_to_shared(&smem_var));

    size_t smem_ptr_64bit = static_cast<size_t>(smem_ptr_32bit); // zero-extend to 64 bits
    void*  generic_ptr    = __cvta_shared_to_generic(smem_ptr_64bit);
    assert(generic_ptr == &smem_var);

    T __ldg(const T* address);

    T __ldcg(const T* address);
    T __ldca(const T* address);
    T __ldcs(const T* address);
    T __ldlu(const T* address);
    T __ldcv(const T* address);

    void __stwb(T* address, T value);
    void __stcg(T* address, T value);
    void __stcs(T* address, T value);
    void __stwt(T* address, T value);

    void __trap();

    __device__ void __nanosleep(unsigned nanoseconds);

    __device__ void mutex_lock(unsigned* mutex) {
        unsigned ns = 8;
        while (atomicCAS(mutex, 0, 1) == 1) {
            __nanosleep(ns);
            if (ns < 256) {
                ns *= 2;
            }
        }
    }

    __device__ void mutex_unlock(unsigned *mutex) {
        atomicExch(mutex, 0);
    }

         int __vimax3_s32  (     int,      int,      int);
    unsigned __vimax3_s16x2(unsigned, unsigned, unsigned);
    unsigned __vimax3_u32  (unsigned, unsigned, unsigned);
    unsigned __vimax3_u16x2(unsigned, unsigned, unsigned);

         int __vimin3_s32  (     int,      int,      int);
    unsigned __vimin3_s16x2(unsigned, unsigned, unsigned);
    unsigned __vimin3_u32  (unsigned, unsigned, unsigned);
    unsigned __vimin3_u16x2(unsigned, unsigned, unsigned);

         int __vimax_s32_relu  (     int,      int);
    unsigned __vimax_s16x2_relu(unsigned, unsigned);

         int __vimin_s32_relu  (     int,      int);
    unsigned __vimin_s16x2_relu(unsigned, unsigned);

         int __vimax3_s32_relu  (     int,      int,      int);
    unsigned __vimax3_s16x2_relu(unsigned, unsigned, unsigned);

         int __vimin3_s32_relu  (     int,      int,      int);
    unsigned __vimin3_s16x2_relu(unsigned, unsigned, unsigned);

         int __vibmax_s32  (     int,      int, bool* pred);
    unsigned __vibmax_u32  (unsigned, unsigned, bool* pred);
    unsigned __vibmax_s16x2(unsigned, unsigned, bool* pred);
    unsigned __vibmax_u16x2(unsigned, unsigned, bool* pred);

         int __vibmin_s32  (     int,      int, bool* pred);
    unsigned __vibmin_u32  (unsigned, unsigned, bool* pred);
    unsigned __vibmin_s16x2(unsigned, unsigned, bool* pred);
    unsigned __vibmin_u16x2(unsigned, unsigned, bool* pred);

         int __viaddmax_s32  (     int,     int,       int);
    unsigned __viaddmax_s16x2(unsigned, unsigned, unsigned);
    unsigned __viaddmax_u32  (unsigned, unsigned, unsigned);
    unsigned __viaddmax_u16x2(unsigned, unsigned, unsigned);

         int __viaddmin_s32  (     int,     int,       int);
    unsigned __viaddmin_s16x2(unsigned, unsigned, unsigned);
    unsigned __viaddmin_u32  (unsigned, unsigned, unsigned);
    unsigned __viaddmin_u16x2(unsigned, unsigned, unsigned);

         int __viaddmax_s32_relu  (     int,      int,      int);
    unsigned __viaddmax_s16x2_relu(unsigned, unsigned, unsigned);

         int __viaddmin_s32_relu  (     int,      int,      int);
    unsigned __viaddmin_s16x2_relu(unsigned, unsigned, unsigned);

    int a           = -15;
    int b           = 8;
    int c           = 5;
    int max_value_0 = __vimax3_s32_relu(a, b, c); // max(-15, 8, 5, 0) = 8
    int d           = -2;
    int e           = -4;
    int max_value_1 = __vimax3_s32_relu(a, d, e); // max(-15, -2, -4, 0) = 0

    int a           = -5;
    int b           = 6;
    int c           = -2;
    int max_value_0 = __viaddmax_s32_relu(a, b, c); // max(-5 + 6, -2, 0) = max(1, -2, 0) = 1
    int d           = 4;
    int max_value_1 = __viaddmax_s32_relu(a, d, c); // max(-5 + 4, -2, 0) = max(-1, -2, 0) = 0

    unsigned a = 9;
    unsigned b = 6;
    bool     smaller_value;
    unsigned min_value = __vibmin_u32(a, b, &smaller_value); // min_value is 6, smaller_value is true

    unsigned a         = 0x00050002;
    unsigned b         = 0x00070004;
    unsigned c         = 0x00020006;
    unsigned max_value = __vimax3_u16x2(a, b, c); // max(5, 7, 2) and max(2, 4, 6), so max_value is 0x00070006

    struct MyStruct {
        static constexpr int value = 4;
    };

    inline constexpr int Count = 4;

    __device__ void foo(int* p1, int* p2) {
        // no argument specified, the loop will be completely unrolled
        #pragma unroll
        for (int i = 0; i < 12; ++i)
            p1[i] += p2[i] * 2;

        // unroll value = 5
        #pragma unroll (Count + 1)
        for (int i = 0; i < 12; ++i)
            p1[i] += p2[i] * 4;

        // unroll value = 1, loop unrolling disabled
        #pragma unroll 1
        for (int i = 0; i < 12; ++i)
            p1[i] += p2[i] * 8;

        // unroll value = 4
        #pragma unroll (MyStruct::value)
        for (int i = 0; i < 12; ++i)
            p1[i] += p2[i] * 16;

        // negative value, pragma unroll ignored
        #pragma unroll -1
        for (int i = 0; i < 12; ++i)
            p1[i] += p2[i] * 2;
    }

    void* __builtin_assume_aligned(const void* ptr, size_t align)
    void* __builtin_assume_aligned(const void* ptr, size_t align, <integral type> offset)

    void* res1 = __builtin_assume_aligned(ptr, 32);    // compiler can assume 'res1' is at least 32-byte aligned
    void* res2 = __builtin_assume_aligned(ptr, 32, 8); // compiler can assume 'res2 = (char*) ptr - 8' is at least 32-byte aligned

    void __builtin_assume(bool predicate)
    void __assume        (bool predicate) // only with Microsoft Compiler

    __device__ bool is_greater_than_zero(int value) {
        return value > 0;
    }

    __device__ bool f(int value) {
        __builtin_assume(value > 0);
        return is_greater_than_zero(value); // returns true, without evaluating the condition
    }

    long __builtin_expect(long input, long expected)

    // indicate to the compiler that likely "var == 0"
    if (__builtin_expect(var, 0))
        doit();

    void __builtin_unreachable(void)

    // indicates to the compiler that the default case label is never reached.
    switch (in) {
        case 1:  return 4;
        case 2:  return 10;
        default: __builtin_unreachable();
    }

    #pragma nv_abi preserve_n_data(EXPR) preserve_n_control(EXPR)

    #pragma nv_abi preserve_n_data(16)
    __device__ void dev_func();

    #pragma nv_abi preserve_n_data(16) preserve_n_control(8)
    __device__ int dev_func() {
        return 0;
    }

    __device__ int dev_func1();

    struct MyStruct {
        __device__ int member_func2();
    };

    __device__ void test() {
        auto* dev_func_ptr = &dev_func1; // type: int (*)(void)
        #pragma nv_abi preserve_n_control(8)
        int v1 = dev_func_ptr();         // CORRECT, indirect call

        #pragma nv_abi preserve_n_control(8)
        int v2 = dev_func1();            // WRONG, direct call; the pragma has no effect
                                         // dev_func1 has type: int(void)

        auto& dev_func_ref = &dev_func1; // type: int (&)(void)
        #pragma nv_abi preserve_n_control(8)
        int v3 = dev_func_ref();         // WRONG, call to a reference
                                         // the pragma has no effect

        auto member_function_ptr = &MyStruct::member_func2; // type: int (MyStruct::*)(void)
        #pragma nv_abi preserve_n_control(8)
        int v4 = member_function_ptr();  // WRONG, indirect call to member function
                                         // the pragma has no effect
    }

    #pragma nv_abi preserve_n_control(8)
    __device__ int dev_func3();

    __device__ int dev_func4();

    __device__ void test() {
        int v1 = dev_func3();            // CORRECT, the pragma affects the direct call

        auto* dev_func_ptr = &dev_func4; // type: int (*)(void)
        #pragma nv_abi preserve_n_control(8)
        int v2 = dev_func_ptr();         // CORRECT, the pragma affects the indirect call

        int v3 = dev_func_ptr();         // WRONG, the pragma has no effect
    }

    #pragma nv_mma_throughput
    __global__  void kernel(){
    ...
    }

    #define assert(expression) /* unspecified */

    #include <assert.h>

     __global__ void testAssert(void) {
         int is_one        = 1;
         int should_be_one = 0;

         // This will have no effect
         assert(is_one);

         // This will halt kernel execution
         assert(should_be_one);
     }

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

    void __brkpt();

    #pragma nv_diag_suppress
    #pragma nv_diag_warning
    #pragma nv_diag_error
    #pragma nv_diag_default
    #pragma nv_diag_once

    #pragma nv_diag_xxx <error_number1>, <error_number2> ...

    #pragma nv_diag_suppress 177 // "declared but never referenced"
    void foo() {
        int i = 0;
    }

    #pragma nv_diag_default 177
    void bar() {
        int i = 0;
    }

    #pragma nv_diagnostic push
    #pragma nv_diagnostic pop

    #pragma nv_diagnostic push
    #pragma nv_diag_suppress 177 // "declared but never referenced"
    void foo() {
        int i = 0;
    }

    #pragma nv_diagnostic pop
    void bar() {
        int i = 0; // raise a warning
    }

    template<typename Use, int m, int n, int k, typename T, typename Layout=void> class fragment;

    void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned ldm);
    void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned ldm, layout_t layout);
    void store_matrix_sync(T* mptr, const fragment<...> &a, unsigned ldm, layout_t layout);
    void fill_fragment(fragment<...> &a, const T& v);
    void mma_sync(fragment<...> &d, const fragment<...> &a, const fragment<...> &b, const fragment<...> &c, bool satf=false);

    enum fragment<Use, m, n, k, T, Layout>::num_elements;
    T fragment<Use, m, n, k, T, Layout>::x[num_elements];

    wmma::fragment<wmma::accumulator, 16, 16, 16, float> frag;
    float alpha = 0.5f; // Same value for all threads in warp
    /*...*/
    for(int t=0; t<frag.num_elements; t++)
    frag.x[t] *= alpha;

    precision::tf32 -> float

    namespace experimental {
        namespace precision {
            struct u4; // 4-bit unsigned
            struct s4; // 4-bit signed
            struct b1; // 1-bit
       }
        enum bmmaBitOp {
            bmmaBitOpXOR = 1, // compute_75 minimum
            bmmaBitOpAND = 2  // compute_80 minimum
        };
        enum bmmaAccumulateOp { bmmaAccumulateOpPOPC = 1 };
    }

    experimental::precision::u4 -> unsigned (8 elements in 1 storage element)
    experimental::precision::s4 -> int (8 elements in 1 storage element)
    experimental::precision::b1 -> unsigned (32 elements in 1 storage element)
    T -> T  //all other types

    fragA.cu: void foo() { wmma::fragment<...> mat_a; bar(&mat_a); }
    fragB.cu: void bar(wmma::fragment<...> *mat_a) { // operate on mat_a }

    // sm_70 fragment layout
    $> nvcc -dc -arch=compute_70 -code=sm_70 fragA.cu -o fragA.o
    // sm_75 fragment layout
    $> nvcc -dc -arch=compute_75 -code=sm_75 fragB.cu -o fragB.o
    // Linking the two together
    $> nvcc -dlink -arch=sm_75 fragA.o fragB.o -o frag.o

    #include <mma.h>
    using namespace nvcuda;

    __global__ void wmma_ker(half *a, half *b, float *c) {
       // Declare the fragments
       wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::col_major> a_frag;
       wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> b_frag;
       wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

       // Initialize the output to zero
       wmma::fill_fragment(c_frag, 0.0f);

       // Load the inputs
       wmma::load_matrix_sync(a_frag, a, 16);
       wmma::load_matrix_sync(b_frag, b, 16);

       // Perform the matrix multiplication
       wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

       // Store the output
       wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
    }