### Run Quickstart Example

Source: https://docs.nvidia.com/cuda/cutile-python/quickstart.html

Execute the quickstart vector addition example from the command line. This verifies the setup and installation.

```bash
$ python3 samples/quickstart/VectorAdd_quickstart.py
```

--------------------------------

### Install CuPy for cuTile Examples

Source: https://docs.nvidia.com/cuda/cutile-python/quickstart.html

Install the CuPy package, which is used by some cuTile Python samples, including the quickstart example.

```bash
pip install cupy-cuda13x
```

--------------------------------

### cuSPARSE csrilu02() Example Setup

Source: https://docs.nvidia.com/cuda/cusparse/index.html

This example demonstrates the initial setup for using csrilu02() to solve a preconditioned system M*y = x, where M is an incomplete LU factorization of A. It includes creating descriptors and setting matrix properties.

```c
// Suppose that A is m x m sparse matrix represented by CSR format,
// Assumption:
// - handle is already created by cusparseCreate(),
// - (d_csrRowPtr, d_csrColInd, d_csrVal) is CSR of A on device memory,
// - d_x is right hand side vector on device memory,
// - d_y is solution vector on device memory.
// - d_z is intermediate result on device memory.

cusparseMatDescr_t descr_M = 0;
cusparseMatDescr_t descr_L = 0;
cusparseMatDescr_t descr_U = 0;
csrilu02Info_t info_M  = 0;
csrsv2Info_t  info_L  = 0;
csrsv2Info_t  info_U  = 0;
int pBufferSize_M;
int pBufferSize_L;
int pBufferSize_U;
int pBufferSize;
void *pBuffer = 0;
int structural_zero;
int numerical_zero;
const double alpha = 1.;
const cusparseSolvePolicy_t policy_M = CUSPARSE_SOLVE_POLICY_NO_LEVEL;
const cusparseSolvePolicy_t policy_L = CUSPARSE_SOLVE_POLICY_NO_LEVEL;
const cusparseSolvePolicy_t policy_U = CUSPARSE_SOLVE_POLICY_USE_LEVEL;
const cusparseOperation_t trans_L  = CUSPARSE_OPERATION_NON_TRANSPOSE;
const cusparseOperation_t trans_U  = CUSPARSE_OPERATION_NON_TRANSPOSE;

// step 1: create a descriptor which contains
// - matrix M is base-1
// - matrix L is base-1
// - matrix L is lower triangular
// - matrix L has unit diagonal
// - matrix U is base-1
// - matrix U is upper triangular
// - matrix U has non-unit diagonal
cusparseCreateMatDescr(&descr_M);
cusparseSetMatIndexBase(descr_M, CUSPARSE_INDEX_BASE_ONE);
```

--------------------------------

### Heterogeneous Batched Memory Transfer Setup

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advanced/advanced-host-programming.html

Illustrates the setup for a heterogeneous batched memory transfer, where different transfers within the batch might have varying characteristics. This example shows allocating buffers for a portion of the batch.

```C++
std::vector<void *> srcs(batch_size);
std::vector<void *> dsts(batch_size);
std::vector<size_t> sizes(batch_size);

// Allocate the src and dst buffers
for (size_t i = 0; i < batch_size - 10; i++) {
    cudaMallocHost(&srcs[i], sizes[i]);
    cudaMalloc(&dsts[i], sizes[i]);
}

int buffer[10];

for (size_t i = batch_size - 10; i < batch_size; i++) {

```

--------------------------------

### Run cuTile Python Samples

Source: https://docs.nvidia.com/cuda/cutile-python/quickstart.html

Invoke other cuTile Python samples directly from the command line, similar to how the quickstart example is run.

```bash
$ python3 samples/FFT.py
```

--------------------------------

### cuBLASDx Introduction Example

Source: https://docs.nvidia.com/cuda/cublasdx/using_cublasdx.html

Demonstrates the setup and execution of GEMM operations using cuBLASDx with managed memory allocation. This example showcases both shared memory and register fragment APIs for GEMM.

```C++
template<unsigned int Arch>
int introduction_example() {
    using GEMM = decltype(cublasdx::Size<32, 32, 32>()
                  + cublasdx::Precision<double>()
                  + cublasdx::Type<cublasdx::type::real>()
                  + cublasdx::Arrangement<cublasdx::row_major, cublasdx::col_major>()
                  + cublasdx::Function<cublasdx::function::MM>()
                  + cublasdx::SM<890>()
                  + cublasdx::Block()
                  + cublasdx::BlockDim<256>());

    using value_type = typename example::uniform_value_type_t<GEMM>;

    constexpr auto global_a_size = example::global_memory_size_of<GEMM>::a_size;
    constexpr auto global_b_size = example::global_memory_size_of<GEMM>::b_size;
    constexpr auto global_c_size = example::global_memory_size_of<GEMM>::c_size;

    // Allocate managed memory for A, B, C matrices in one go
    value_type* abc;
    auto        size       = global_a_size + global_b_size + global_c_size;
    auto        size_bytes = size * sizeof(value_type);
    CUDA_CHECK_AND_EXIT(cudaMallocManaged(&abc, size_bytes));
    // Generate data
    for (size_t i = 0; i < size; i++) {
        abc[i] = double(i / size);
    }

    value_type* a = abc;
    value_type* b = abc + global_a_size;
    value_type* c = abc + global_a_size + global_b_size;


    // Shared memory API: C = alpha * A * B + beta * C
    // Invokes kernel with GEMM::block_dim threads in CUDA block
    gemm_kernel_shared<GEMM><<<1, GEMM::block_dim, cublasdx::get_shared_storage_size<GEMM>()>>>(1.0, a, b, 0.5, c);

    // Register fragment API: C = alpha * A * B + beta * C
    // Invokes kernel with GEMM::block_dim threads in CUDA block
    gemm_kernel_registers<GEMM><<<1, GEMM::block_dim, cublasdx::get_shared_storage_size_ab<GEMM>()>>>(a, b, c);

    CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
    CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());

    CUDA_CHECK_AND_EXIT(cudaFree(abc));
    std::cout << "Success" << std::endl;
    return 0;
}
```

--------------------------------

### CURAND Poisson Distribution Example Setup

Source: https://docs.nvidia.com/cuda/curand/device-api-overview.html

Includes necessary headers, macros for CUDA and CURAND calls, constants for simulation, and an enum for API types. This setup is required for the Poisson distribution simulation.

```c++
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <curand_kernel.h>
#include <curand.h>

#define CUDA_CALL(x) do { if((x) != cudaSuccess) { 
    printf("Error at %s:%d\n",__FILE__,__LINE__); 
    return EXIT_FAILURE;}} while(0)
#define CURAND_CALL(x) do { if((x)!=CURAND_STATUS_SUCCESS) { 
    printf("Error at %s:%d\n",__FILE__,__LINE__);
    return EXIT_FAILURE;}} while(0)


#define HOURS 16
#define OPENING_HOUR 7
#define CLOSING_HOUR (OPENING_HOUR + HOURS)

#define access_2D(type, ptr, row, column, pitch)
    *((type*)((char*)ptr + (row) * pitch) + column)

enum API_TYPE {
    HOST_API = 0,
    SIMPLE_DEVICE_API = 1,
    ROBUST_DEVICE_API = 2,
};

/* global variables */
API_TYPE api;
int report_break;
int cashiers_load_h[HOURS];
__constant__ int cashiers_load[HOURS];
```

--------------------------------

### mbarrier.arrive_drop Example with Register

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

A basic example demonstrating the declaration of a 32-bit register `cnt` which can be used with `mbarrier.arrive_drop` operations. This setup is typical for managing synchronization counts in PTX.

```assembly
.reg .b32 cnt;

```

--------------------------------

### Example of solving a precondition system with incomplete Cholesky factorization

Source: https://docs.nvidia.com/cuda/cusparse/index.html

This example demonstrates solving the precondition system M*y = x, where M is the product of a Cholesky factorization L and its transpose (M=LLH). It outlines the necessary setup and variable declarations for using cusparse functions.

```c
// Suppose that A is m x m sparse matrix represented by CSR format,
// Assumption:
// - handle is already created by cusparseCreate(),
// - (d_csrRowPtr, d_csrColInd, d_csrVal) is CSR of A on device memory,
// - d_x is right hand side vector on device memory,
// - d_y is solution vector on device memory.
// - d_z is intermediate result on device memory.

cusparseMatDescr_t descr_M = 0;
cusparseMatDescr_t descr_L = 0;
csric02Info_t info_M  = 0;
csrsv2Info_t  info_L  = 0;
csrsv2Info_t  info_Lt = 0;
int pBufferSize_M;
int pBufferSize_L;
int pBufferSize_Lt;
int pBufferSize;
void *pBuffer = 0;
int structural_zero;
int numerical_zero;
const double alpha = 1.;
const cusparseSolvePolicy_t policy_M  = CUSPARSE_SOLVE_POLICY_NO_LEVEL;
const cusparseSolvePolicy_t policy_L  = CUSPARSE_SOLVE_POLICY_NO_LEVEL;
```

--------------------------------

### Download and Install Keyring (Ubuntu Network Install)

Source: https://docs.nvidia.com/cuda/cuda-installation-guide-linux

Downloads the NVIDIA CUDA repository keyring and installs it for network cross repository installation on Ubuntu.

```bash
$ wget https://developer.download.nvidia.com/compute/cuda/repos/<distro>/cross-linux-<arch>/cuda-keyring_1.1-1_all.deb
# dpkg -i cuda-keyring_1.1-1_all.deb

```

--------------------------------

### Download and install CUDA keyring (WSL - Network Repo)

Source: https://docs.nvidia.com/cuda/cuda-quick-start-guide/index.html

Download and install the CUDA keyring package for network repository installations in WSL.

```bash
wget https://developer.download.nvidia.com/compute/cuda/repos/<distro>/<arch>/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb

```

--------------------------------

### Install CUDA using RPM on OpenSUSE

Source: https://docs.nvidia.com/cuda/cuda-quick-start-guide

Installs the CUDA repository, refreshes the package cache, and installs the CUDA toolkit using Zypper.

```bash
sudo rpm --install cuda-repo-<distro>-<version>.<architecture>.rpm
sudo rpm --erase gpg-pubkey-7fa2af80*
sudo zypper refresh
sudo zypper install cuda
```

--------------------------------

### Example with Libraries, Host Linker, and Dynamic Parallelism

Source: https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html

Demonstrates a complex build process involving device linking, static library creation, and host linking with multiple libraries.

```bash
nvcc --gpu-architecture=sm_100 --device-c a.cu b.cu
vnvcc --gpu-architecture=sm_100 --device-link a.o b.o --output-file link.o
vnvcc --lib --output-file libgpu.a a.o b.o link.o
g++ host.o --library=gpu --library-path=<path> \
    --library=cudadevrt --library=cudart
```

--------------------------------

### Introduction Example: Basic FFT Launch

Source: https://docs.nvidia.com/cuda/cufftdx/introduction1.html

Demonstrates launching a basic FFT kernel. It sets up the FFT description and launches the kernel with specified block dimensions and shared memory size. Assumes CUDA error checking macros are defined.

```cpp
#include <cufftdx.hpp>
using namespace cufftdx;

template<class FFT>
__global__ void block_fft_kernel(FFT::value_type* data,
                                  typename FFT::workspace_type workspace) {
  using complex_type = typename FFT::value_type;
  // Registers
  complex_type thread_data[FFT::storage_size];

  // FFT::shared_memory_size bytes of shared memory
  extern __shared__ __align__(alignof(float4))
    complex_type shared_mem[];

  // Execute FFT
  FFT().execute(thread_data, shared_mem, workspace);
}

// CUDA_CHECK_AND_EXIT - marco checks if function returns
// cudaSuccess; if not it prints the error code and exits the
// program
void introduction_example(FFT::value_type* data
    /* data is a manage memory pointer*/) {
  // FFT description
  using FFT = decltype(Size<128>() + Precision<float>() + Type<fft_type::c2c>()
                       + Direction<fft_direction::forward>() + FFTsPerBlock<1>()
                       + ElementsPerThread<8>() + SM<750>() + Block());

  cudaError_t error_code = cudaSuccess;
  auto workspace = make_workspace<FFT>(error_code);
  CUDA_CHECK_AND_EXIT(error_code);

  // Invokes kernel with FFT::block_dim threads in CUDA block
  block_fft_kernel<FFT><<<1, FFT::block_dim,
                          FFT::shared_memory_size>>>(data, workspace);
  CUDA_CHECK_AND_EXIT(cudaPeekAtLastError());
  CUDA_CHECK_AND_EXIT(cudaDeviceSynchronize());
}
```

--------------------------------

### Floating Point Add Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution

Provides concrete examples of using the 'add' instruction with specific modifiers.

```assembly
@p  add.rz.ftz.f32  f1,f2,f3;
add.rp.ftz.f32x2    d, a, b;

```

--------------------------------

### Example: Running deviceQuery on Orin Platforms

Source: https://docs.nvidia.com/cuda/cuda-for-tegra-appnote/index.html

This example demonstrates how to set LD_LIBRARY_PATH and run the deviceQuery sample application on Orin platforms.

```bash
LD_LIBRARY_PATH=/usr/local/cuda-13.3/compat_orin:$LD_LIBRARY_PATH ~/Samples/1_Utilities/deviceQuery

```

--------------------------------

### NVRTC Bundled Headers Example (C++)

Source: https://docs.nvidia.com/cuda/nvrtc/index.html

Demonstrates installing and using NVRTC's bundled headers to compile a kernel with half-precision types. This example requires CUDA 13.3+ and installs headers to a version-specific directory.

```cpp
#include <nvrtc.h>
#include <iostream>
#include <string>

#define NVRTC_SAFE_CALL(x)                                        \
  do {                                                            \
    nvrtcResult result = x;                                       \
    if (result != NVRTC_SUCCESS) {                                \
      std::cerr << "\nerror: " #x " failed with error "           \
                << nvrtcGetErrorString(result) << '\n';           \
      exit(1);                                                    \
    }                                                             \
  } while(0)

const char *hsaxpy =                                                \
"  #include <cuda_fp16.h>                                       \n" \
"  extern \"C\" __global__                                      \n" \
"  void hsaxpy(__half a, __half *x, __half *y,                  \n" \
"              __half *out, size_t n)                            \n" \
"  {                                                            \n" \
"    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;        \n" \
"    if (tid < n) {                                             \n" \
"      out[tid] = a * x[tid] + y[tid];                          \n" \
"    }                                                          \n" \
"  }                                                            \n";

int main()
{
  // Query bundled headers version and build a version-specific install path.
  nvrtcBundledHeadersInfo info;
  nvrtcGetBundledHeadersInfo(&info, nullptr);
  std::string installPath = "cuda_headers_"
      + std::to_string(info.cudaVersionMajor) + "."
      + std::to_string(info.cudaVersionMinor);

  // Install bundled headers.
  const char* errorLog = nullptr;
  nvrtcResult res = nvrtcInstallBundledHeaders(
      installPath.c_str(), NVRTC_INSTALL_HEADERS_SKIP_IF_EXISTS, &errorLog);
  if (res != NVRTC_SUCCESS) {
    std::cerr << "Install failed: "
              << (errorLog ? errorLog : nvrtcGetErrorString(res)) << '\n';
    return 1;
  }

  // Create and compile the program with bundled headers.
  nvrtcProgram prog;
  NVRTC_SAFE_CALL(
    nvrtcCreateProgram(&prog, hsaxpy, "hsaxpy.cu", 0, NULL, NULL));

  std::string includeOpt = "-I" + installPath;
  const char *opts[] = { includeOpt.c_str() };
  nvrtcResult compileResult = nvrtcCompileProgram(prog, 1, opts);

  // Obtain compilation log from the program.
  size_t logSize;
  NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
  char *log = new char[logSize];
  NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
  std::cout << log << '\n';
  delete[] log;
  if (compileResult != NVRTC_SUCCESS) {
    exit(1);
  }

  // Obtain PTX from the program.
  size_t ptxSize;
  NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &ptxSize));
  char *ptx = new char[ptxSize];
  NVRTC_SAFE_CALL(nvrtcGetPTX(prog, ptx));
  NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));

  std::cout << "Successfully compiled hsaxpy kernel ("
            << ptxSize << " bytes of PTX)\n";

  delete[] ptx;
  return 0;
}

```

--------------------------------

### cuFFTDx Introduction Example: Setting up and Launching FFT Kernel

Source: https://docs.nvidia.com/cuda/cufftdx/introduction2.html

This example demonstrates how to set up a cuFFTDx FFT execution descriptor, allocate managed memory, generate input data, create a workspace, and launch the `block_fft_kernel`. It uses suggested `FFTsPerBlock` for optimal performance with default `ElementsPerThread`.

```C++
#include <cufftdx.hpp>
using namespace cufftdx;

template<class FFT>
__global__ void block_fft_kernel(FFT::value_type* data,
                                  typename FFT::workspace_type workspace) {
  using complex_type = typename FFT::value_type;

  // Registers
  complex_type thread_data[FFT::storage_size];

  // Local batch id of this FFT in CUDA block, in range
  // [0; FFT::ffts_per_block)
  const unsigned int local_fft_id = threadIdx.y;
  // Global batch id of this FFT in CUDA grid is equal to number
  // of batches per CUDA block (ffts_per_block) times CUDA block
  // id, plus local batch id.
  const unsigned int global_fft_id =
      (blockIdx.x * FFT::ffts_per_block) + local_fft_id;

  // Load data from global memory to registers
  const unsigned int offset = cufftdx::size_of<FFT>::value * global_fft_id;
  const unsigned int stride = FFT::stride;
  unsigned int       index  = offset + threadIdx.x;
  for (unsigned int i = 0; i < FFT::elements_per_thread; i++) {
      // Make sure not to go out-of-bounds
      if ((i * stride + threadIdx.x) < cufftdx::size_of<FFT>::value) {
          thread_data[i] = data[index];
          index += stride;
      }
  }

  // FFT::shared_memory_size bytes of shared memory
  extern __shared__ __align__(alignof(float4)) complex_type shared_mem[];

  // Execute FFT
  FFT().execute(thread_data, shared_mem, workspace);

  // Save results
  index = offset + threadIdx.x;
  for (unsigned int i = 0; i < FFT::elements_per_thread; i++) {
      if ((i * stride + threadIdx.x) < cufftdx::size_of<FFT>::value) {
          data[index] = thread_data[i];
          index += stride;
      }
  }
}

void introduction_example(cudaStream_t stream = 0) {
  // Base of the FFT description
  using FFT_base = decltype(Size<128>() + Precision<float>() + 
      Type<fft_type::c2c>() + 
      Direction<fft_direction::forward>() 
      /* Notice lack of ElementsPerThread and FFTsPerBlock 
       * operators */ 
      + SM<750>() + Block());
  // FFT description with suggested FFTs per CUDA block for the
  // default (optimal) elements per thread
  using FFT = decltype(FFT_base() + 
      FFTsPerBlock<FFT_base::suggested_ffts_per_block>());

  // Allocate managed memory for input/output
  complex_type* data;
  auto          size       = FFT::ffts_per_block * cufftdx::size_of<FFT>::value;
  auto          size_bytes = size * sizeof(complex_type);
  cudaMallocManaged(&data, size_bytes);
  // Generate data
  for (size_t i = 0; i < size; i++) {
      data[i] = complex_type {float(i), -float(i)};
  }

  cudaError_t error_code = cudaSuccess;


  auto workspace = make_workspace<FFT>(error_code, stream);

  // Invokes kernel with FFT::block_dim threads in CUDA block
  block_fft_kernel<FFT><<<1, FFT::block_dim, 
                          FFT::shared_memory_size,
                          stream>>>(data, workspace);
  cudaDeviceSynchronize();

  cudaFree(data);
}

```

--------------------------------

### Install CUDA Toolkit

Source: https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html

Installs the CUDA Toolkit using the DNF package manager after repository setup.

```bash
# dnf install cuda-toolkit

```

--------------------------------

### PTX Sign-extend or Zero-extend (szext) Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Provides examples of the szext instruction, showcasing both .clamp and .wrap modes for sign and zero extension. The second example demonstrates a specific case where the result is 0.

```assembly
szext.clamp.s32 rd, ra, rb;
szext.wrap.u32  rd, 0xffffffff, 0; // Result is 0.
```

--------------------------------

### TRSM Execution Example

Source: https://docs.nvidia.com/cuda/cublasdx/api/trsm_methods.html

An example demonstrating the setup and execution of a TRSM operation using shared memory and the tensor API.

```APIDOC
## TRSM Execution Example

### Description
This example demonstrates how to set up shared memory and execute a TRSM operation using the cuBLASDx library. It involves copying data to shared memory, executing the TRSM operation, and copying the result back to global memory.

### Code Example
```cpp
using TRSM = decltype(cublasdx::Size<32, 16>()
                      + cublasdx::Precision<float>()
                      + cublasdx::Type<cublasdx::type::real>()
                      + cublasdx::Function<cublasdx::function::TRSM>()
                      + cublasdx::Side<cublasdx::side::left>()
                      + cublasdx::FillMode<cublasdx::fill_mode::lower>()
                      + cublasdx::Diag<cublasdx::diag::non_unit>()
                      + cublasdx::SM<900>()
                      + cublasdx::Block());

extern __shared__ __align__(16) cublasdx::byte smem[];

using alignment = cublasdx::alignment_of<TRSM>;

// Get batches for global memory A and B
auto batch_a = cublasdx::get_batch(global_a, TRSM::get_layout_gmem_a(), blockIdx.x);
auto batch_b = cublasdx::get_batch(global_b, TRSM::get_layout_gmem_b(), blockIdx.x);

// Slice shared memory for A and B with appropriate alignments and layouts
auto [smem_a, smem_b] = cublasdx::shared_memory::slice<float, float>(
    smem,
    alignment::a, TRSM::get_layout_smem_a(),
    alignment::b, TRSM::get_layout_smem_b());

// Copy data from global memory to shared memory
cublasdx::copy<TRSM, alignment::a>(batch_a, smem_a);
cublasdx::copy<TRSM, alignment::b>(batch_b, smem_b);
cublasdx::copy_wait();

// Execute the TRSM operation in shared memory
TRSM{}.execute(smem_a, smem_b);
__syncthreads();

// Copy the result from shared memory back to global memory
cublasdx::copy<TRSM, alignment::b>(smem_b, batch_b);
```

### Notes
- This example assumes `global_a` and `global_b` are pointers to the global memory matrices.
- `cublasdx::copy_wait()` is used to ensure all copy operations are complete before proceeding.
- `__syncthreads()` is crucial after the `execute` call to synchronize threads before reading the results from `B`.
- For a complete example, refer to example 17 in the cuBLASDx Examples.
```

--------------------------------

### Fabric try_red Source and Completion Mechanism Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Provides example configurations for the source memory space and completion mechanisms used with fabric.try_red.

```syntax
.src = { .shared::cta }
```

```syntax
.completion_mechanism0 = { .mbarrier::complete_tx::16B.mbarrier::report::fabric }
```

```syntax
.completion_mechanism1 = { .mbarrier::complete_tx::16B.mbarrier::report::fabric.counted::bytes }
```

--------------------------------

### vmad Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Provides concrete examples of using the vmad instruction with saturation and scaling.

```assembly
vmad.s32.s32.u32.sat    r0, r1, r2, -r3;
vmad.u32.u32.u32.shr15  r0, r1.h0, r2.h0, r3;


```

--------------------------------

### Launch and Replace Hints for cuTile Kernel

Source: https://docs.nvidia.com/cuda/cutile-python/execution.html

Demonstrates launching a cuTile kernel with specific occupancy hints and then recompiling it with updated hints using `replace_hints`. This is useful for tuning kernel performance by adjusting compilation parameters.

```python
import cuda.tile as ct
import torch

torch.cuda.init()
stream = torch.cuda.current_stream()

@ct.kernel(occupancy=2)
def kernel():
    pass

# compile
ct.launch(torch.cuda.current_stream(), (1,), kernel, ())
# cache hit
ct.launch(torch.cuda.current_stream(), (1,), kernel, ())

new_kernel = kernel.replace_hints(occupancy=4)

# compile with new hints
ct.launch(torch.cuda.current_stream(), (1,), new_kernel, ())
# cache hit
ct.launch(torch.cuda.current_stream(), (1,), new_kernel, ())
```

--------------------------------

### Get Default Global and Shared Memory Layouts

Source: https://docs.nvidia.com/cuda/cublasdx/api/other_methods.html

Use these functions to get the default memory layouts for matrices A, B, or C in global or shared memory. No setup is required.

```cpp
__forceinline__ __host__ __device__ constexpr static auto get_layout_gmem_a();
__forceinline__ __host__ __device__ constexpr static auto get_layout_gmem_b();
__forceinline__ __host__ __device__ constexpr static auto get_layout_gmem_c();

__forceinline__ __host__ __device__  constexpr static auto get_layout_smem_a();
__forceinline__ __host__ __device__  constexpr static auto get_layout_smem_b();
__forceinline__ __host__ __device__  constexpr static auto get_layout_smem_c();
```

--------------------------------

### CUDA cp.reduce.async.bulk.tensor Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Provides concrete examples of using the cp.reduce.async.bulk.tensor instruction with different dimensions, reduction operations, load modes, and completion mechanisms.

```cuda-ptx
cp.reduce.async.bulk.tensor.1d.global.shared::cta.add.tile.bulk_group
                                             [tensorMap0, {tc0}], [sMem0];


```

```cuda-ptx
cp.reduce.async.bulk.tensor.2d.global.shared::cta.and.bulk_group.L2::cache_hint
                                             [tensorMap1, {tc0, tc1}], [sMem1] , policy;


```

```cuda-ptx
cp.reduce.async.bulk.tensor.3d.global.shared::cta.xor.im2col.bulk_group
                                             [tensorMap2, {tc0, tc1, tc2}], [sMem2];


```

--------------------------------

### Fabric try_red Semantic and Scope Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Illustrates example settings for the semantic (relaxed) and scope (system) parameters in fabric.try_red operations.

```syntax
.sem = { .relaxed }
```

```syntax
.scope = { .sys }
```

--------------------------------

### Example: Launching Kernel with Cooperative Attribute

Source: https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html

Demonstrates how to set the cooperative attribute for a kernel launch using CUlaunchConfig and CUlaunchAttribute.

```c
CUlaunchAttribute coopAttr = {.id = CU_LAUNCH_ATTRIBUTE_COOPERATIVE,
                                     .value = 1};
       CUlaunchConfig config = {... // set block and grid dimensions
                              .attrs = &coopAttr,
                              .numAttrs = 1};
      
       cuLaunchKernelEx(&config, kernel, NULL, NULL);
```

--------------------------------

### Install Bundled Headers with Versioned Path

Source: https://docs.nvidia.com/cuda/nvrtc/index.html

This snippet shows how to get NVRTC bundled header information and construct a version-specific installation path to prevent using stale headers after an NVRTC update.

```c++
nvrtcBundledHeadersInfo info;
vrtcGetBundledHeadersInfo(&info, nullptr);

std::string installPath = "/app/data/nvrtc_headers_"
    + std::to_string(info.cudaVersionMajor) + "."
    + std::to_string(info.cudaVersionMinor);

vrtcInstallBundledHeaders(installPath.c_str(),
                           NVRTC_INSTALL_HEADERS_SKIP_IF_EXISTS, nullptr);


```

--------------------------------

### selp Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Provides examples of using the selp instruction for selecting between 32-bit signed integers and 32-bit floating-point values.

```assembly
    selp.s32  r0,r,g,p;
@q  selp.f32  f0,t,x,xp;


```

--------------------------------

### min Instruction Examples

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Provides examples of using the 'min' instruction with different data types and modifiers.

```assembly
    min.s32  r0,a,b;
@p  min.u16  h,i,j;
    min.s16x2.relu u,v,w;
    min.u8x4 p, q, r;

```

--------------------------------

### Create and Initialize JIT Linker

Source: https://docs.nvidia.com/cuda/nvjitlink/index.html

Creates an instance of the JIT linker and initializes it with specified options. This is the first step before adding input files.

```c
nvJitLink_t linker;
const char* link_options[] = { "-arch=sm_80" };
vJitLinkCreate(&linker, 1, link_options);
```

--------------------------------

### CURAND Poisson API Example

Source: https://docs.nvidia.com/cuda/curand/device-api-overview.html

This example demonstrates how to use the CURAND device API to generate Poisson-distributed random numbers. It covers setup, generation using different APIs (HOST_API, SIMPLE_DEVICE_API, ROBUST_DEVICE_API), and cleanup.

```c++
curandState *devStates;
unsigned int *devResults, *hostResults;
unsigned int *poisson_numbers_d;
curandDiscreteDistribution_t poisson_1, poisson_2;
curandDiscreteDistribution_t poisson_3;
curandGenerator_t gen;

/* Setting cashiers, report and API */
settings();

/* Allocate space for results on device */
CUDA_CALL(cudaMallocPitch((void **)&devResults, &pitch,
            64 * 64 * sizeof(unsigned int), 60 * HOURS));

/* Allocate space for results on host */
hostResults = (unsigned int *)calloc(pitch * 60 * HOURS,
            sizeof(unsigned int));

/* Allocate space for prng states on device */
CUDA_CALL(cudaMalloc((void **)&devStates, 64 * 64 *
              sizeof(curandState)));

/* Setup prng states */
if (api != HOST_API){
    setup_kernel<<<64, 64>>>(devStates);
}
/* Simulate queue  */
switch (api){
    case HOST_API:
        /* Create pseudo-random number generator */
        CURAND_CALL(curandCreateGenerator(&gen,
                            CURAND_RNG_PSEUDO_DEFAULT));
        /* Set seed */
        CURAND_CALL(curandSetPseudoRandomGeneratorSeed(
                                            gen, 1234ULL));
        /* compute n */
        n = 64 * 64 * HOURS * 60;
        /* Allocate n unsigned ints on device */
        CUDA_CALL(cudaMalloc((void **)&poisson_numbers_d,
                            n * sizeof(unsigned int)));
        /* Generate n unsigned ints on device */
        CURAND_CALL(curandGeneratePoisson(gen,
                            poisson_numbers_d, n, 4.0));
        host_API_kernel<<<64, 64>>>(poisson_numbers_d,
                                        devResults, pitch);
        /* Cleanup */
        CURAND_CALL(curandDestroyGenerator(gen));
        break;
    case SIMPLE_DEVICE_API:
        simple_device_API_kernel<<<64, 64>>>(devStates,
                                        devResults, pitch);
        break;
    case ROBUST_DEVICE_API:
        /* Create histograms for Poisson(1) */
        CURAND_CALL(curandCreatePoissonDistribution(1.0,
                                                &poisson_1));
        /* Create histograms for Poisson(2) */
        CURAND_CALL(curandCreatePoissonDistribution(2.0,
                                                &poisson_2));
        /* Create histograms for Poisson(3) */
        CURAND_CALL(curandCreatePoissonDistribution(3.0,
                                                &poisson_3));
        robust_device_API_kernel<<<64, 64>>>(devStates,
                        poisson_1, poisson_2, poisson_3,
                        devResults, pitch);
        /* Cleanup */
        CURAND_CALL(curandDestroyDistribution(poisson_1));
        CURAND_CALL(curandDestroyDistribution(poisson_2));
        CURAND_CALL(curandDestroyDistribution(poisson_3));
        break;
    default:
        fprintf(stderr, "Wrong API\n");
}
/* Copy device memory to host */
CUDA_CALL(cudaMemcpy2D(hostResults, pitch, devResults,
            pitch, 64 * 64 * sizeof(unsigned int),
            60 * HOURS, cudaMemcpyDeviceToHost));
/* Show result */
print_statistics(hostResults, pitch);
/* Cleanup */
CUDA_CALL(cudaFree(devStates));
CUDA_CALL(cudaFree(devResults));
free(hostResults);
return EXIT_SUCCESS;
}

```

--------------------------------

### Compile and Run deviceQuery Sample

Source: https://docs.nvidia.com/cuda/cuda-installation-guide-microsoft-windows

This command is used to build and run the deviceQuery sample program to verify hardware and software configuration. Assumes default installation directory structure.

```bash
deviceQuery
```

--------------------------------

### Example PTX Version Directives

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Examples of valid .version directives in PTX modules, indicating different PTX ISA versions.

```plaintext
.version 3.1
```

```plaintext
.version 3.0
```

```plaintext
.version 2.3
```

--------------------------------

### Dynamic Parallelism Hello World Example

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/dynamic-parallelism.html

A basic 'Hello World' program demonstrating dynamic parallelism, where a parent kernel launches a child kernel and a tail kernel into a special stream.

```cuda
#include <stdio.h>

__global__ void childKernel()
{
    printf("Hello ");
}

__global__ void tailKernel()
{
    printf("World!\n");
}

__global__ void parentKernel()
{
    // launch child
    childKernel<<<1,1>>>();
    if (cudaSuccess != cudaGetLastError()) {
        return;
    }

    // launch tail into cudaStreamTailLaunch stream
    // implicitly synchronizes: waits for child to complete
    tailKernel<<<1,1,0,cudaStreamTailLaunch>>>();

}

int main(int argc, char *argv[])
{
    // launch parent
    parentKernel<<<1,1>>>();
    if (cudaSuccess != cudaGetLastError()) {
        return 1;
    }

    // wait for parent to complete
    if (cudaSuccess != cudaDeviceSynchronize()) {
        return 2;
    }

    return 0;
}


```

--------------------------------

### Extents Deduction Guide Example

Source: https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/extents.html

Demonstrates the use of the deduction guide for ct::extents, enabling CTAD (Class Template Argument Deduction) from integral or ct::integral_constant arguments. This simplifies the declaration of extents objects.

```cpp
namespace ct = ::cuda::tiles;
using namespace ct::literals;
ct::extents x{4_ic, 7};
```

--------------------------------

### CUDA API Forward Progress Guarantee Example 3

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cuda-cpp-execution-model.html

Similar to Example 2, this shows that a single CUDA query API call does not guarantee device thread start, potentially leading to host thread unresponsiveness.

```cuda-cpp
// Example: Execution.Model.API.3
// Allowed outcome: eventually, no thread makes progress.
// Rationale: same as Example.Model.API.2, with the addition that a single CUDA query API call does not guarantee
// the device thread eventually starts, only repeated CUDA query API calls do (see Execution.Model.API.4).
cuda::atomic<int, cuda::thread_scope_system> flag = 0;
__global__ void producer() { flag.store(1); }
int main() {
    cudaHostRegister(&flag, sizeof(flag));
    producer<<<1,1>>>();
    (void)cudaStreamQuery(0);
    while (flag.load() == 0);
    return cudaDeviceSynchronize();
}
```

--------------------------------

### NVCC Example with Libraries and Dynamic Parallelism

Source: https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc

Demonstrates a compilation flow involving device linking, static library creation, and host linking with CUDA runtime and device runtime libraries.

```bash
nvcc --gpu-architecture=sm_100 --device-c a.cu b.cu
vncc --gpu-architecture=sm_100 --device-link a.o b.o --output-file link.o
vncc --lib --output-file libgpu.a a.o b.o link.o
g++ host.o --library=gpu --library-path=<path> \
    --library=cudadevrt --library=cudart
```

--------------------------------

### Runtime Fatbinary Creation Example

Source: https://docs.nvidia.com/cuda/nvfatbin/contents.html

Demonstrates how to create a fatbinary at runtime. This example is provided as a C++ source file.

```cpp
#include <iostream>
#include <vector>
#include <cuda.h>

// Dummy function to represent some CUDA kernel logic
__global__ void myKernel() {
    // Kernel code goes here
}

int main() {
    CUresult res;
    CUcontext context;
    CUdevice device;
    CUmodule module;
    CUfunction function;

    // Initialize CUDA
    res = cuInit(0);
    if (res != CUDA_SUCCESS) {
        std::cerr << "cuInit failed: " << res << std::endl;
        return 1;
    }

    // Get device and create context
    res = cuDeviceGet(&device, 0);
    if (res != CUDA_SUCCESS) {
        std::cerr << "cuDeviceGet failed: " << res << std::endl;
        return 1;
    }

    res = cuCtxCreate(&context, 0, device);
    if (res != CUDA_SUCCESS) {
        std::cerr << "cuCtxCreate failed: " << res << std::endl;
        return 1;
    }

    // In a real scenario, you would load PTX or CUBIN code here.
    // For demonstration, we'll simulate module loading.
    // This part is highly simplified and won't actually create a functional fatbinary.
    std::cout << "Simulating fatbinary creation..." << std::endl;

    // In a real application, you would use cuModuleLoadData or similar
    // to load compiled PTX or CUBIN code into a module.
    // For example:
    // const char *ptx_code = 
    //    "\"version=5.0\"\n... PTX code ...";
    // res = cuModuleLoadDataEx(&module, ptx_code, CU_JIT_TARGET_FROM_ENVIRONMENT);

    // Placeholder for module and function retrieval
    // module = nullptr; // Replace with actual loaded module
    // res = cuModuleGetFunction(&function, module, "myKernel");

    std::cout << "Fatbinary creation simulation complete." << std::endl;

    // Clean up CUDA context
    res = cuCtxDestroy(context);
    if (res != CUDA_SUCCESS) {
        std::cerr << "cuCtxDestroy failed: " << res << std::endl;
        return 1;
    }

    return 0;
}

```

--------------------------------

### Install Third-party Libraries (OpenSUSE)

Source: https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html

Installs necessary third-party development libraries for CUDA samples on OpenSUSE using zypper.

```bash
# zypper install freeglut-devel libX11-devel libXi-devel libXmu-devel make Mesa-libGL-devel freeimage-devel
```

--------------------------------

### Get cuBLAS Logger Callback

Source: https://docs.nvidia.com/cuda/cublas

Retrieves the function pointer to a previously installed custom user-defined callback function for cuBLAS logging.

```c
cublasStatus_t cublasGetLoggerCallback(
    cublasLogCallback* userCallback)


```

--------------------------------

### Install Third-party Libraries (OpenSUSE)

Source: https://docs.nvidia.com/cuda/cuda-installation-guide-linux

Installs necessary third-party development libraries for CUDA samples on OpenSUSE.

```bash
# zypper install freeglut-devel libX11-devel libXi-devel libXmu-devel make Mesa-libGL-devel freeimage-devel

```

--------------------------------

### PTX Pragma: frequency Example

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Specifies the execution frequency hint for a basic block. This pragma should be placed at the start of the basic block.

```cuda-ptx
.entry foo (...)
{
    .pragma "frequency 32";
    ...
}
```

--------------------------------

### CUDA Graph Setup for Sibling Launch

Source: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html

Demonstrates the setup process for launching a CUDA graph using the sibling launch mechanism. This involves creating, instantiating, and uploading a device graph, then creating and instantiating a host graph that launches the device graph.

```cuda
void graphSetup() {
    cudaGraphExec_t gExec1, gExec2;
    cudaGraph_t g1, g2;

    // Create, instantiate, and upload the device graph.
    create_graph(&g2);
    cudaGraphInstantiate(&gExec2, g2, cudaGraphInstantiateFlagDeviceLaunch);
    cudaGraphUpload(gExec2, stream);

    // Create and instantiate the launching graph.
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
    launchSiblingGraph<<<1, 1, 0, stream>>>(gExec2);
    cudaStreamEndCapture(stream, &g1);
    cudaGraphInstantiate(&gExec1, g1);

    // Launch the host graph, which will in turn launch the device graph.
    cudaGraphLaunch(gExec1, stream);

}

```

--------------------------------

### Example LTO Helper Execution

Source: https://docs.nvidia.com/cuda/cufftdx/lto_helper.html

An example command demonstrating how to run the LTO Helper with a specific output folder, input CSV file, and CUDA architectures.

```bash
./cufftdx_cufft_lto_helper output_folder input.csv --CUDA_ARCHITECTURES=75

```

--------------------------------

### Example: Copying Matrix Between Global and Shared Memory

Source: https://docs.nvidia.com/cuda/cublasdx/api/other_tensors.html

This example demonstrates copying a matrix 'A' from global memory to shared memory and then back to global memory using `cublasdx::copy` and `cublasdx::copy_wait`. It includes setup for shared memory and tensor definitions.

```cpp
using BLAS = decltype(Size<128, 128, 128>() + Type<type::real>() + Precision<float, float, double>() + Block() + ...);
extern __shared__ __align__(16) cublasdx::byte smem[];

// Slice shared memory
auto [smem_a, smem_b, smem_c] = cublasdx::slice_shared_memory<BLAS>(smem);

auto gmem_tensor_a = cublasdx::make_tensor(a_gmem_pointer, BLAS::get_layout_gmem_a());
auto smem_tensor_a = cublasdx::make_tensor(smem_a, BLAS::suggest_layout_smem_a());

// Copy from global to shared
using alignment = cublasdx::alignment_of<BLAS>;
cublasdx::copy<BLAS, alignment::a>(gmem_tensor_a, smem_tensor_a);
cublasdx::copy_wait();

// Copy from shared to global
cublasdx::copy<BLAS, alignment::a>(smem_tensor_a, gmem_tensor_a);
cublasdx::copy_wait();

```

--------------------------------

### Compile and Run cuFFTMp Example

Source: https://docs.nvidia.com/cuda/cufftmp/getting_started/index.html

This command shows how to compile and run the cuFFTMp example. Ensure that the include and library paths for cuFFT, nvshmem, and MPI are correctly specified.

```bash
$ nvcc app.cu -I/path/to/cufftMp/include -L/path/to/cufftMp/lib -lcufftMp  -L/path/to/nvshmem/lib -lnvshmem_host -I/path/to/mpi/include -L/path/to/mpi/lib -lmpi -o app
$ mpirun -n 2 ./app
PASSED

```

--------------------------------

### nvjpegCreateSimple()

Source: https://docs.nvidia.com/cuda/nvjpeg/index.html

Allocates and initializes the library handle with default codec implementations and memory allocators. This is a simplified way to get started with the nvjpeg library.

```APIDOC
## nvjpegCreateSimple()

### Description
Allocates and initializes the library handle, with default codec implementations selected by library and default memory allocators.

### Signature
```c
nvjpegStatus_t nvjpegCreateSimple(nvjpegHandle_t *handle);
```

### Parameters
#### Input/Output Parameters
- **handle** (nvjpegHandle_t *) - Host - The library handle.

### Returns
`nvjpegStatus_t` - An error code as specified in nvJPEG API Return Codes.
```

--------------------------------

### Atomic Add System Example

Source: https://docs.nvidia.com/cuda/cuda-c-programming-guide

Demonstrates atomic addition to a memory location from both CPU and GPU using `atomicAdd_system` and `__sync_fetch_and_add`.

```cuda-c
__global__ void mykernel(int *addr) {
  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
}

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

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

```

--------------------------------

### CUDBGAPI_st::getGridDim

Source: https://docs.nvidia.com/cuda/debugger-api/group__GRID.html

Get the dimensions in blocks of the given grid. This function retrieves the grid's dimensions in blocks, with support starting from CUDA 4.0.

```APIDOC
## CUDBGAPI_st::getGridDim

### Description
Get the dimensions in blocks of the given grid. Since CUDA 4.0.

### Method
Not specified (likely a function call in an SDK)

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Parameters
- **dev** (uint32_t) - device index
- **sm** (uint32_t) - SM index
- **wp** (uint32_t) - warp index
- **gridDim** (CuDim3*) - the dimensions of the grid

### Returns
CUDBG_SUCCESS, CUDBG_ERROR_INVALID_ARGS, CUDBG_ERROR_UNINITIALIZED, CUDBG_ERROR_INVALID_WARP, CUDBG_ERROR_INITIALIZATION_FAILURE, CUDBG_ERROR_INVALID_ATTRIBUTE, CUDBG_ERROR_INVALID_CONTEXT, CUDBG_ERROR_RECURSIVE_API_CALL
```