NVIDIA CUDA Toolkit

https://docs.nvidia.com/cuda/
Admin
The NVIDIA CUDA Toolkit provides a development environment for creating high-performance...
Tokens:24,985,918
Snippets:122,312
Trust Score:10
Update:18 hours ago
Show doc for...
Tokens:
Raw
### Full example: includes, error handling, and tensor mode/extent setup

Source: https://docs.nvidia.com/cuda/cutensor/2.2.0/getting_started.html

Complete main function from the cuTENSOR getting started example. Includes necessary headers, defines error handling macros, sets up data types, and creates mode and extent vectors for tensors A, B, and C. This is the starting point for the contraction example.

```cpp
#include <stdlib.h>
#include <stdio.h>
#include <assert.h>

#include <cuda_runtime.h>
#include <cutensor.h>

#include <unordered_map>
#include <vector>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x)                                             \
{ const auto err = x;                                               \
    if( err != CUTENSOR_STATUS_SUCCESS )                              \
    { printf("Error: %s\n", cutensorGetErrorString(err)); exit(-1); } \
};

#define HANDLE_CUDA_ERROR(x)                                      \
{ const auto err = x;                                             \
    if( err != cudaSuccess )                                        \
    { printf("Error: %s\n", cudaGetErrorString(err)); exit(-1); } \
};

int main(int argc, char** argv)
{
    // Host element type definition
    typedef float floatTypeA;
    typedef float floatTypeB;
    typedef float floatTypeC;
    typedef float floatTypeCompute;

    // CUDA types
    cutensorDataType_t typeA = CUTENSOR_R_32F;
    cutensorDataType_t typeB = CUTENSOR_R_32F;
    cutensorDataType_t typeC = CUTENSOR_R_32F;
    cutensorComputeDescriptor_t descCompute = CUTENSOR_COMPUTE_DESC_32F;

    printf("Include headers and define data types\n");

    /* ***************************** */

    // Create vector of modes
    std::vector<int> modeC{'m','u','n','v'};
    std::vector<int> modeA{'m','h','k','n'};
    std::vector<int> modeB{'u','k','v','h'};
    int nmodeA = modeA.size();
    int nmodeB = modeB.size();
    int nmodeC = modeC.size();

    // Extents
    std::unordered_map<int, int64_t> extent;
    extent['m'] = 96;
    extent['n'] = 96;
    extent['u'] = 96;
    extent['v'] = 64;
    extent['h'] = 64;
    extent['k'] = 64;

    // Create a vector of extents for each tensor
    std::vector<int64_t> extentC;
    for(auto mode : modeC)
        extentC.push_back(extent[mode]);
    std::vector<int64_t> extentA;
    for(auto mode : modeA)
        extentA.push_back(extent[mode]);
    std::vector<int64_t> extentB;
    for(auto mode : modeB)
        extentB.push_back(extent[mode]);
```

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

### Full Example with Plan Cache

Source: https://docs.nvidia.com/cuda/cutensor/1.7.0/plan_cache.html

Complete introductory example that enhances the Getting Started example with plan cache usage. It includes attaching cachelines, a loop around the contraction call, and detaching cachelines. The code defines tensor modes and extents, allocates device memory, and uses the HANDLE_ERROR macro for error checking.

```cpp
  1#include <stdlib.h>
  2#include <stdio.h>
  3
  4#include <cuda_runtime.h>
  5#include <cutensor.h>
  6
  7#include <unordered_map>
  8#include <vector>
  9
 10// Handle cuTENSOR errors
 11#define HANDLE_ERROR(x) {                                                              \
 12  const auto err = x;                                                                  \
 13  if( err != CUTENSOR_STATUS_SUCCESS )                                                   \
 14  { printf("Error: %s in line %d\n", cutensorGetErrorString(err), __LINE__); exit(-1); } \
 15}
 16
 17int main(int argc, char** argv)
 18{
 19  // Host element type definition
 20  typedef float floatTypeA;
 21  typedef float floatTypeB;
 22  typedef float floatTypeC;
 23  typedef float floatTypeCompute;
 24
 25  // CUDA types
 26  cudaDataType_t typeA = CUDA_R_32F;
 27  cudaDataType_t typeB = CUDA_R_32F;
 28  cudaDataType_t typeC = CUDA_R_32F;
 29  cutensorComputeType_t typeCompute = CUTENSOR_COMPUTE_32F;
 30
 31  floatTypeCompute alpha = (floatTypeCompute)1.1f;
 32  floatTypeCompute beta  = (floatTypeCompute)0.9f;
 33
 34  printf("Include headers and define data types\n");
 35
 36  /* ***************************** */
 37
 38  // Create vector of modes
 39  std::vector<int> modeC{'m','u','n','v'};
 40  std::vector<int> modeA{'m','h','k','n'};
 41  std::vector<int> modeB{'u','k','v','h'};
 42  int nmodeA = modeA.size();
 43  int nmodeB = modeB.size();
 44  int nmodeC = modeC.size();
 45
 46  // Extents
 47  std::unordered_map<int, int64_t> extent;
 48  extent['m'] = 96;
 49  extent['n'] = 96;
 50  extent['u'] = 96;
 51  extent['v'] = 64;
 52  extent['h'] = 64;
 53  extent['k'] = 64;
 54
 55  // Create a vector of extents for each tensor
 56  std::vector<int64_t> extentC;
 57  for(auto mode : modeC)
 58      extentC.push_back(extent[mode]);
 59  std::vector<int64_t> extentA;
 60  for(auto mode : modeA)
 61      extentA.push_back(extent[mode]);
 62  std::vector<int64_t> extentB;
 63  for(auto mode : modeB)
 64      extentB.push_back(extent[mode]);
 65
 66  printf("Define modes and extents\n");
 67
 68  /* ***************************** */
 69
 70  // Number of elements of each tensor
 71  size_t elementsA = 1;
 72  for(auto mode : modeA)
 73      elementsA *= extent[mode];
 74  size_t elementsB = 1;
 75  for(auto mode : modeB)
 76      elementsB *= extent[mode];
 77  size_t elementsC = 1;
 78  for(auto mode : modeC)
 79      elementsC *= extent[mode];
 80
 81  // Size in bytes
 82  size_t sizeA = sizeof(floatTypeA) * elementsA;
 83  size_t sizeB = sizeof(floatTypeB) * elementsB;
 84  size_t sizeC = sizeof(floatTypeC) * elementsC;
 85
 86  // Allocate on device
 87  void *A_d, *B_d, *C_d;
 88  cudaMalloc((void**)&A_d, sizeA);
```

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

### Complete cuTENSOR example: setup, data allocation, and tensor descriptors

Source: https://docs.nvidia.com/cuda/cutensor/2.0.1/getting_started.html

Full C++ example that includes headers, error handling macros, main function with data type definitions, mode and extent setup, device/host allocation, data initialization, and cuTENSOR handle creation. The example is cut off at the tensor descriptor creation section.

```cpp
#include <stdlib.h>
#include <stdio.h>
#include <assert.h>

#include <cuda_runtime.h>
#include <cutensor.h>

#include <unordered_map>
#include <vector>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x)                                             \
{ const auto err = x;                                               \
  if( err != CUTENSOR_STATUS_SUCCESS )                              \
  { printf("Error: %s\n", cutensorGetErrorString(err)); exit(-1); } \
};

#define HANDLE_CUDA_ERROR(x)                                      \
{ const auto err = x;                                             \
  if( err != cudaSuccess )                                        \
  { printf("Error: %s\n", cudaGetErrorString(err)); exit(-1); } \
};

int main(int argc, char** argv)
{
  // Host element type definition
  typedef float floatTypeA;
  typedef float floatTypeB;
  typedef float floatTypeC;
  typedef float floatTypeCompute;

  // CUDA types
  cutensorDataType_t typeA = CUTENSOR_R_32F;
  cutensorDataType_t typeB = CUTENSOR_R_32F;
  cutensorDataType_t typeC = CUTENSOR_R_32F;
  cutensorComputeDescriptor_t descCompute = CUTENSOR_COMPUTE_DESC_32F;

  printf("Include headers and define data types\n");

  /* ***************************** */

  // Create vector of modes
  std::vector<int> modeC{'m','u','n','v'};
  std::vector<int> modeA{'m','h','k','n'};
  std::vector<int> modeB{'u','k','v','h'};
  int nmodeA = modeA.size();
  int nmodeB = modeB.size();
  int nmodeC = modeC.size();

  // Extents
  std::unordered_map<int, int64_t> extent;
  extent['m'] = 96;
  extent['n'] = 96;
  extent['u'] = 96;
  extent['v'] = 64;
  extent['h'] = 64;
  extent['k'] = 64;

  // Create a vector of extents for each tensor
  std::vector<int64_t> extentC;
  for(auto mode : modeC)
      extentC.push_back(extent[mode]);
  std::vector<int64_t> extentA;
  for(auto mode : modeA)
      extentA.push_back(extent[mode]);
  std::vector<int64_t> extentB;
  for(auto mode : modeB)
      extentB.push_back(extent[mode]);

  printf("Define modes and extents\n");

  /* ***************************** */

  // Number of elements of each tensor
  size_t elementsA = 1;
  for(auto mode : modeA)
      elementsA *= extent[mode];
  size_t elementsB = 1;
  for(auto mode : modeB)
      elementsB *= extent[mode];
  size_t elementsC = 1;
  for(auto mode : modeC)
      elementsC *= extent[mode];

  // Size in bytes
  size_t sizeA = sizeof(floatTypeA) * elementsA;
  size_t sizeB = sizeof(floatTypeB) * elementsB;
  size_t sizeC = sizeof(floatTypeC) * elementsC;

  // Allocate on device
  void *A_d, *B_d, *C_d;
  cudaMalloc((void**)&A_d, sizeA);
  cudaMalloc((void**)&B_d, sizeB);
  cudaMalloc((void**)&C_d, sizeC);

  // Allocate on host
  floatTypeA *A = (floatTypeA*) malloc(sizeof(floatTypeA) * elementsA);
  floatTypeB *B = (floatTypeB*) malloc(sizeof(floatTypeB) * elementsB);
  floatTypeC *C = (floatTypeC*) malloc(sizeof(floatTypeC) * elementsC);

  // Initialize data on host
  for(int64_t i = 0; i < elementsA; i++)
      A[i] = (((float) rand())/RAND_MAX - 0.5)*100;
  for(int64_t i = 0; i < elementsB; i++)
      B[i] = (((float) rand())/RAND_MAX - 0.5)*100;
  for(int64_t i = 0; i < elementsC; i++)
      C[i] = (((float) rand())/RAND_MAX - 0.5)*100;

  // Copy to device
  HANDLE_CUDA_ERROR(cudaMemcpy(A_d, A, sizeA, cudaMemcpyHostToDevice));
  HANDLE_CUDA_ERROR(cudaMemcpy(B_d, B, sizeB, cudaMemcpyHostToDevice));
  HANDLE_CUDA_ERROR(cudaMemcpy(C_d, C, sizeC, cudaMemcpyHostToDevice));

  const uint32_t kAlignment = 128; // Alignment of the global-memory device pointers (bytes)
  assert(uintptr_t(A_d) % kAlignment == 0);
  assert(uintptr_t(B_d) % kAlignment == 0);
  assert(uintptr_t(C_d) % kAlignment == 0);

  printf("Allocate, initialize and transfer tensors\n");

  /*************************
   * cuTENSOR
   *************************/

  cutensorHandle_t handle;
  HANDLE_ERROR(cutensorCreate(&handle));

  /**********************
   * Create Tensor Descriptors
   **********************/

```

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

### Run the quickstart vector addition sample

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

Command to run the quickstart sample from the command line. If setup is correct, it prints '✓ vector_add_example passed!'.

```bash
$ python3 samples/quickstart/VectorAdd_quickstart.py
✓ vector_add_example passed!
```

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

### Source file for G1: introduction_example.cu

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

Path to the source file for the descriptor and launch example.

```text
example/cublasdx/01_gemm_introduction/introduction_example.cu
```

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

### Complete Example with Plan Cache

Source: https://docs.nvidia.com/cuda/cutensor/1.6.2/plan_cache.html

A full cuTENSOR example that includes the plan cache setup and a loop around the contraction call. It defines tensor modes and extents, allocates device memory, and uses the cache to improve performance. The example is based on the Getting Started example enhanced with cache-related calls.

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

#include <cuda_runtime.h>
#include <cutensor.h>

#include <unordered_map>
#include <vector>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x) {                                                              \
  const auto err = x;                                                                  \
  if( err != CUTENSOR_STATUS_SUCCESS )                                                   \
  { printf("Error: %s in line %d\n", cutensorGetErrorString(err), __LINE__); exit(-1); } \
}

int main(int argc, char** argv)
{
  // Host element type definition
  typedef float floatTypeA;
  typedef float floatTypeB;
  typedef float floatTypeC;
  typedef float floatTypeCompute;

  // CUDA types
  cudaDataType_t typeA = CUDA_R_32F;
  cudaDataType_t typeB = CUDA_R_32F;
  cudaDataType_t typeC = CUDA_R_32F;
  cutensorComputeType_t typeCompute = CUTENSOR_COMPUTE_32F;

  floatTypeCompute alpha = (floatTypeCompute)1.1f;
  floatTypeCompute beta  = (floatTypeCompute)0.9f;

  printf("Include headers and define data types\n");

  /* ***************************** */

  // Create vector of modes
  std::vector<int> modeC{'m','u','n','v'};
  std::vector<int> modeA{'m','h','k','n'};
  std::vector<int> modeB{'u','k','v','h'};
  int nmodeA = modeA.size();
  int nmodeB = modeB.size();
  int nmodeC = modeC.size();

  // Extents
  std::unordered_map<int, int64_t> extent;
  extent['m'] = 96;
  extent['n'] = 96;
  extent['u'] = 96;
  extent['v'] = 64;
  extent['h'] = 64;
  extent['k'] = 64;

  // Create a vector of extents for each tensor
  std::vector<int64_t> extentC;
  for(auto mode : modeC)
      extentC.push_back(extent[mode]);
  std::vector<int64_t> extentA;
  for(auto mode : modeA)
      extentA.push_back(extent[mode]);
  std::vector<int64_t> extentB;
  for(auto mode : modeB)
      extentB.push_back(extent[mode]);

  printf("Define modes and extents\n");

  /* ***************************** */

  // Number of elements of each tensor
  size_t elementsA = 1;
  for(auto mode : modeA)
      elementsA *= extent[mode];
  size_t elementsB = 1;
  for(auto mode : modeB)
      elementsB *= extent[mode];
  size_t elementsC = 1;
  for(auto mode : modeC)
      elementsC *= extent[mode];

  // Size in bytes
  size_t sizeA = sizeof(floatTypeA) * elementsA;
  size_t sizeB = sizeof(floatTypeB) * elementsB;
  size_t sizeC = sizeof(floatTypeC) * elementsC;

  // Allocate on device
  void *A_d, *B_d, *C_d;
  cudaMalloc((void**)&A_d, sizeA);
```

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

### Example with libraries, host linker, and dynamic parallelism

Source: https://docs.nvidia.com/cuda/pdf/CUDA_Compiler_Driver_NVCC.pdf

Commands for an example using libraries, host linker, and dynamic parallelism.

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

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

### Distributed Matmul with NVFP4 - Setup Example

Source: https://docs.nvidia.com/cuda/nvmath-python/latest/distributed-apis/linalg/generated/nvmath.distributed.linalg.advanced.matmul.html

Example code showing how to initialize the nvmath.distributed context, get the process group and rank, and create input tensors for a distributed matmul operation.

```APIDOC
## Distributed Matmul with NVFP4 - Setup Example

### Description
Example code showing how to initialize the nvmath.distributed context, get the process group and rank, and create input tensors for a distributed matmul operation.

### Code
```python
>>> import cupy as cp
>>> import nvmath.distributed
>>> from nvmath.distributed.distribution import Slab
```

Get process group used to initialize nvmath.distributed (for information on initializing `nvmath.distributed`, you can refer to the documentation or to the Matmul examples in [nvmath/examples/distributed/linalg/advanced](https://github.com/NVIDIA/nvmath-python/tree/main/examples/distributed/linalg/advanced/matmul)):

```python
>>> process_group = nvmath.distributed.get_context().process_group
```

Get my process rank:

```python
>>> rank = process_group.rank
```

Create three float32 ndarrays on the GPU:

```python
>>> # (code continues)
```
```

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

### Show help for qiskit_ghz.py example

Source: https://docs.nvidia.com/cuda/cuquantum/latest/getting-started/index.html

Displays the help message for the qiskit_ghz.py example script inside the cuQuantum Appliance container.

```bash
docker exec cuquantum-appliance python /home/cuquantum/examples/qiskit_ghz.py --help
```

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

### Introductory JIT compilation example: full C++ program

Source: https://docs.nvidia.com/cuda/cutensor/2.8.0/just_in_time_compilation.html

This is the complete C++ program for the introductory JIT compilation example. It computes a tensor contraction with many modes, using single-precision complex values and 3XTF32 arithmetic. The code includes error handling macros, CPU and GPU timers, and the main function that sets up the contraction by defining modes, extents, and allocating data. It is the same steps as the Getting Started guide but with a different contraction example to emphasize the benefit of JIT compilation when the number of contracted modes increases.

```cpp
#include <chrono>
#include <complex>
#include <stdlib.h>
#include <stdio.h>
#include <unordered_map>
#include <vector>

#include <cuda_runtime.h>
#include <cutensor.h>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x) {                                                                \
  const auto err = x;                                                                    \
  if( err != CUTENSOR_STATUS_SUCCESS )                                                   \
  { printf("Error: %s in line %d\n", cutensorGetErrorString(err), __LINE__); exit(-1); } \
};

// Handle CUDA errors
#define HANDLE_CUDA_ERROR(x) {                                                       \
  const auto err = x;                                                                \
  if( err != cudaSuccess )                                                           \
  { printf("Error: %s in line %d\n", cudaGetErrorString(err), __LINE__); exit(-1); } \
};

class CPUTimer
{
public:
    void start()
    {
        start_ = std::chrono::steady_clock::now();
    }

    double seconds()
    {
        end_ = std::chrono::steady_clock::now();
        elapsed_ = end_ - start_;
        //return in ms
        return elapsed_.count() * 1000;
    }

private:
    typedef std::chrono::steady_clock::time_point tp;
    tp start_;
    tp end_;
    std::chrono::duration<double> elapsed_;
};

struct GPUTimer
{
  GPUTimer()
  {
    cudaEventCreate(&start_);
    cudaEventCreate(&stop_);
    cudaEventRecord(start_, 0);
  }

  ~GPUTimer()
  {
    cudaEventDestroy(start_);
    cudaEventDestroy(stop_);
  }

  void start()
  {
    cudaEventRecord(start_, 0);
  }

  float seconds()
  {
    cudaEventRecord(stop_, 0);
    cudaEventSynchronize(stop_);
    float time;
    cudaEventElapsedTime(&time, start_, stop_);
    return time * 1e-3;
  }
  private:
  cudaEvent_t start_, stop_;
};

int main()
{
  typedef std::complex<float> TypeA;
  typedef std::complex<float> TypeB;
  typedef std::complex<float> TypeC;
  typedef std::complex<float> TypeScalar;

  auto alpha = TypeScalar(1.1, 0.0);
  auto beta  = TypeScalar(0.0, 0.0);

  // CUDA types
  cutensorDataType_t typeA = CUTENSOR_C_32F;
  cutensorDataType_t typeB = CUTENSOR_C_32F;
  cutensorDataType_t typeC = CUTENSOR_C_32F;
  cutensorComputeDescriptor_t descCompute = CUTENSOR_COMPUTE_DESC_3XTF32;


  /* ***************************** */

  // Create vector of modes
  std::vector<int> modeC{0,1,2,3,4,6,8,9,25,26,10,12,14,27,15,28,17,19,29,20,21,30,23,24};
  std::vector<int> modeA{0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16,17,18,19,20,21,22,23,24};
  std::vector<int> modeB{25,26,27,28,29,30,5,7,11,13,16,18,22};
  int nmodeA = modeA.size();
  int nmodeB = modeB.size();
  int nmodeC = modeC.size();

  // Extents
  std::unordered_map<int, int64_t> extent;
  for (auto i = 0; i <= 30; i++)
      extent[i] = 2;

  // Create a vector of extents for each tensor
  std::vector<int64_t> extentC;
  for (auto mode : modeC)
      extentC.push_back(extent[mode]);
  std::vector<int64_t> extentA;
  for (auto mode : modeA)
      extentA.push_back(extent[mode]);
  std::vector<int64_t> extentB;
  for (auto mode : modeB)
      extentB.push_back(extent[mode]);

  /**********************
   * Allocating data
   **********************/

```

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

### Run cuBLASDx examples

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

Run examples directly from their output directories. The examples print Success or Failure for correctness, and timings/throughput/error for performance.

```bash
./build/example/cublasdx/01_gemm_introduction/introduction_example
./build/example/cublasdx/02_gemm_precisions/simple_gemm_fp32
```

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

### cuTENSOR Getting Started Example

Source: https://docs.nvidia.com/cuda/cutensor/2.0.2/getting_started.html

Complete example that initializes cuTENSOR, creates tensor descriptors for tensors A, B, and C, and prepares for a tensor operation. It includes error handling macros and demonstrates memory allocation and data transfer.

```cpp
#include <stdlib.h>
#include <stdio.h>
#include <assert.h>

#include <cuda_runtime.h>
#include <cutensor.h>

#include <unordered_map>
#include <vector>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x)                                             \
{ const auto err = x;                                               \
    if( err != CUTENSOR_STATUS_SUCCESS )                              \
    { printf("Error: %s\n", cutensorGetErrorString(err)); exit(-1); } \
};

#define HANDLE_CUDA_ERROR(x)                                      \
{ const auto err = x;                                             \
    if( err != cudaSuccess )                                        \
    { printf("Error: %s\n", cudaGetErrorString(err)); exit(-1); } \
};

int main(int argc, char** argv)
{
    // Host element type definition
    typedef float floatTypeA;
    typedef float floatTypeB;
    typedef float floatTypeC;
    typedef float floatTypeCompute;

    // CUDA types
    cutensorDataType_t typeA = CUTENSOR_R_32F;
    cutensorDataType_t typeB = CUTENSOR_R_32F;
    cutensorDataType_t typeC = CUTENSOR_R_32F;
    cutensorComputeDescriptor_t descCompute = CUTENSOR_COMPUTE_DESC_32F;

    printf("Include headers and define data types\n");

    /* ***************************** */

    // Create vector of modes
    std::vector<int> modeC{'m','u','n','v'};
    std::vector<int> modeA{'m','h','k','n'};
    std::vector<int> modeB{'u','k','v','h'};
    int nmodeA = modeA.size();
    int nmodeB = modeB.size();
    int nmodeC = modeC.size();

    // Extents
    std::unordered_map<int, int64_t> extent;
    extent['m'] = 96;
    extent['n'] = 96;
    extent['u'] = 96;
    extent['v'] = 64;
    extent['h'] = 64;
    extent['k'] = 64;

    // Create a vector of extents for each tensor
    std::vector<int64_t> extentC;
    for(auto mode : modeC)
        extentC.push_back(extent[mode]);
    std::vector<int64_t> extentA;
    for(auto mode : modeA)
        extentA.push_back(extent[mode]);
    std::vector<int64_t> extentB;
    for(auto mode : modeB)
        extentB.push_back(extent[mode]);

    printf("Define modes and extents\n");

    /* ***************************** */

    // Number of elements of each tensor
    size_t elementsA = 1;
    for(auto mode : modeA)
        elementsA *= extent[mode];
    size_t elementsB = 1;
    for(auto mode : modeB)
        elementsB *= extent[mode];
    size_t elementsC = 1;
    for(auto mode : modeC)
        elementsC *= extent[mode];

    // Size in bytes
    size_t sizeA = sizeof(floatTypeA) * elementsA;
    size_t sizeB = sizeof(floatTypeB) * elementsB;
    size_t sizeC = sizeof(floatTypeC) * elementsC;

    // Allocate on device
    void *A_d, *B_d, *C_d;
    cudaMalloc((void**)&A_d, sizeA);
    cudaMalloc((void**)&B_d, sizeB);
    cudaMalloc((void**)&C_d, sizeC);

    // Allocate on host
    floatTypeA *A = (floatTypeA*) malloc(sizeof(floatTypeA) * elementsA);
    floatTypeB *B = (floatTypeB*) malloc(sizeof(floatTypeB) * elementsB);
    floatTypeC *C = (floatTypeC*) malloc(sizeof(floatTypeC) * elementsC);

    // Initialize data on host
    for(int64_t i = 0; i < elementsA; i++)
        A[i] = (((float) rand())/RAND_MAX - 0.5)*100;
    for(int64_t i = 0; i < elementsB; i++)
        B[i] = (((float) rand())/RAND_MAX - 0.5)*100;
    for(int64_t i = 0; i < elementsC; i++)
        C[i] = (((float) rand())/RAND_MAX - 0.5)*100;

    // Copy to device
    HANDLE_CUDA_ERROR(cudaMemcpy(A_d, A, sizeA, cudaMemcpyHostToDevice));
    HANDLE_CUDA_ERROR(cudaMemcpy(B_d, B, sizeB, cudaMemcpyHostToDevice));
    HANDLE_CUDA_ERROR(cudaMemcpy(C_d, C, sizeC, cudaMemcpyHostToDevice));

    const uint32_t kAlignment = 128; // Alignment of the global-memory device pointers (bytes)
    assert(uintptr_t(A_d) % kAlignment == 0);
    assert(uintptr_t(B_d) % kAlignment == 0);
    assert(uintptr_t(C_d) % kAlignment == 0);

    printf("Allocate, initialize and transfer tensors\n");

    /*************************
     * cuTENSOR
     *************************/ 

    cutensorHandle_t handle;
    HANDLE_ERROR(cutensorCreate(&handle));

    /**********************
     * Create Tensor Descriptors
     **********************/

    cutensorTensorDescriptor_t descA;
    HANDLE_ERROR(cutensorCreateTensorDescriptor(handle,
                &descA,
                nmodeA,
                extentA.data(),
                NULL,/*stride*/
                typeA, kAlignment));

    cutensorTensorDescriptor_t descB;
    HANDLE_ERROR(cutensorCreateTensorDescriptor(handle,
                &descB,
                nmodeB,
                extentB.data(),
                NULL,/*stride*/
                typeB, kAlignment));

    cutensorTensorDescriptor_t descC;
    HANDLE_ERROR(cutensorCreateTensorDescriptor(handle,
                &descC,
                nmodeC,
                extentC.data(),
                NULL,/*stride*/
                typeC, kAlignment));

    printf("Initialize cuTENSOR and tensor descriptors\n");

    return 0;
}
```

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

### Example: Using cuLaunchKernelEx with CUlaunchConfig

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

Example showing how to set up a CUlaunchConfig with a cooperative launch attribute and call cuLaunchKernelEx.

```APIDOC
## Example: Using cuLaunchKernelEx with CUlaunchConfig

### Description
Demonstrates how to configure a kernel launch using CUlaunchConfig and CUlaunchAttribute, then invoke cuLaunchKernelEx.

### Code Example
```
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);
```
```

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

### Full working example with JIT compilation enabled

Source: https://docs.nvidia.com/cuda/cutensor/2.0.0/just_in_time_compilation.html

Complete example that includes the necessary includes, error handling macro, and the JIT-enabled plan preference creation. This is the full code that demonstrates the JIT compilation workflow.

```cpp
#include <chrono>
#include <complex>
#include <stdlib.h>
#include <stdio.h>
#include <unordered_map>
#include <vector>

#include <cuda_runtime.h>
#include <cutensor.h>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x) {                                                                \
  const auto err = x;                                                                    \
  if( err != CUTENSOR_STATUS_SUCCESS )                                                   \
  { printf("Error: %s in line %d\n", cutensorGetErrorString(err), __LINE__); exit(-1); } \
};
```

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

### bfe Example

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html?highlight=proxy

Example usage of the bfe instruction: bfe.b32 d,a,start,len;

```ptx
bfe.b32  d,a,start,len;
```

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

### Complete cuTENSOR initialization and tensor descriptor creation

Source: https://docs.nvidia.com/cuda/cutensor/2.4.1/getting_started.html

This snippet shows the full setup for a cuTENSOR operation: including headers, defining data types, creating tensor descriptors, and allocating memory. It is a complete example that can be used as a starting point for cuTENSOR applications.

```cpp
#include <stdlib.h>
#include <stdio.h>
#include <assert.h>

#include <cuda_runtime.h>
#include <cutensor.h>

#include <unordered_map>
#include <vector>

// Handle cuTENSOR errors
#define HANDLE_ERROR(x)                                             \
{ const auto err = x;                                               \
    if( err != CUTENSOR_STATUS_SUCCESS )                            \
    { printf("Error: %s\n", cutensorGetErrorString(err)); exit(-1); } \
};

#define HANDLE_CUDA_ERROR(x)                                      \
{ const auto err = x;                                             \
    if( err != cudaSuccess )                                        \
    { printf("Error: %s\n", cudaGetErrorString(err)); exit(-1); } \
};

int main(int argc, char** argv)
{
    // Host element type definition
    typedef float floatTypeA;
    typedef float floatTypeB;
    typedef float floatTypeC;
    typedef float floatTypeCompute;

    // CUDA types
    cutensorDataType_t typeA = CUTENSOR_R_32F;
    cutensorDataType_t typeB = CUTENSOR_R_32F;
    cutensorDataType_t typeC = CUTENSOR_R_32F;
    cutensorComputeDescriptor_t descCompute = CUTENSOR_COMPUTE_DESC_32F;

    printf("Include headers and define data types\n");

    /* ***************************** */

    // Create vector of modes
    std::vector<int> modeC{'m','u','n','v'};
    std::vector<int> modeA{'m','h','k','n'};
    std::vector<int> modeB{'u','k','v','h'};
    int nmodeA = modeA.size();
    int nmodeB = modeB.size();
    int nmodeC = modeC.size();

    // Extents
    std::unordered_map<int, int64_t> extent;
    extent['m'] = 96;
    extent['n'] = 96;
    extent['u'] = 96;
    extent['v'] = 64;
    extent['h'] = 64;
    extent['k'] = 64;

    // Create a vector of extents for each tensor
    std::vector<int64_t> extentC;
    for(auto mode : modeC)
        extentC.push_back(extent[mode]);
    std::vector<int64_t> extentA;
    for(auto mode : modeA)
        extentA.push_back(extent[mode]);
    std::vector<int64_t> extentB;
    for(auto mode : modeB)
        extentB.push_back(extent[mode]);

    printf("Define modes and extents\n");

    /* ***************************** */

    // Number of elements of each tensor
    size_t elementsA = 1;
    for(auto mode : modeA)
        elementsA *= extent[mode];
    size_t elementsB = 1;
    for(auto mode : modeB)
        elementsB *= extent[mode];
    size_t elementsC = 1;
    for(auto mode : modeC)
        elementsC *= extent[mode];

    // Size in bytes
    size_t sizeA = sizeof(floatTypeA) * elementsA;
    size_t sizeB = sizeof(floatTypeB) * elementsB;
    size_t sizeC = sizeof(floatTypeC) * elementsC;

    // Allocate on device
    void *A_d, *B_d, *C_d;
    cudaMalloc((void**)&A_d, sizeA);
    cudaMalloc((void**)&B_d, sizeB);
    cudaMalloc((void**)&C_d, sizeC);

    // Allocate on host
    floatTypeA *A = (floatTypeA*) malloc(sizeof(floatTypeA) * elementsA);
    floatTypeB *B = (floatTypeB*) malloc(sizeof(floatTypeB) * elementsB);
    floatTypeC *C = (floatTypeC*) malloc(sizeof(floatTypeC) * elementsC);

    // Initialize data on host
    for(int64_t i = 0; i < elementsA; i++)
        A[i] = (((float) rand())/RAND_MAX - 0.5)*100;
    for(int64_t i = 0; i < elementsB; i++)
        B[i] = (((float) rand())/RAND_MAX - 0.5)*100;
    for(int64_t i = 0; i < elementsC; i++)
        C[i] = (((float) rand())/RAND_MAX - 0.5)*100;

    // Copy to device
    HANDLE_CUDA_ERROR(cudaMemcpy(A_d, A, sizeA, cudaMemcpyHostToDevice));
    HANDLE_CUDA_ERROR(cudaMemcpy(B_d, B, sizeB, cudaMemcpyHostToDevice));
    HANDLE_CUDA_ERROR(cudaMemcpy(C_d, C, sizeC, cudaMemcpyHostToDevice));

    const uint32_t kAlignment = 128; // Alignment of the global-memory device pointers (bytes)
    assert(uintptr_t(A_d) % kAlignment == 0);
    assert(uintptr_t(B_d) % kAlignment == 0);
    assert(uintptr_t(C_d) % kAlignment == 0);

    printf("Allocate, initialize and transfer tensors\n");

    /*************************
     * cuTENSOR
     *************************/ 

    cutensorHandle_t handle;
    HANDLE_ERROR(cutensorCreate(&handle));

    /**********************
     * Create Tensor Descriptors
     **********************/

    cutensorTensorDescriptor_t descA;
    HANDLE_ERROR(cutensorCreateTensorDescriptor(handle,
                &descA,
                nmodeA,
                extentA.data(),
                NULL,/*stride*/
                typeA, kAlignment));

    cutensorTensorDescriptor_t descB;
    HANDLE_ERROR(cutensorCreateTensorDescriptor(handle,
                &descB,
                nmodeB,
                extentB.data(),
                NULL,/*stride*/
                typeB, kAlignment));

    cutensorTensorDescriptor_t descC;
    HANDLE_ERROR(cutensorCreateTensorDescriptor(handle,
                &descC,
                nmodeC,
                extentC.data(),
                NULL,/*stride*/
                typeC, kAlignment));

    printf("Initialize cuTENSOR and tensor descriptors\n");

    return 0;
}
```

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

### Example workflow: build MPO and MPS, create MPO handle

Source: https://docs.nvidia.com/cuda/cuquantum/latest/cudensitymat/examples.html

Shows the beginning of an example workflow: builds the Ising MPO, prints a message if verbose, and creates a cudensitymatMatrixProductOperator handle with open boundary conditions and specified bond dimensions.

```C++
void exampleWorkflow(cudensitymatHandle_t handle)
{
  // --- 1. Build the transverse-field Ising MPO ---
  IsingMPO mpo;
  mpo.build();
  if (verbose)
    std::cout << "Built transverse-field Ising MPO (bond dim " << MPO_BOND_DIM << ")\n";

  const std::vector<int64_t> spaceShape(NUM_SITES, PHYS_DIM);
  std::vector<int64_t> mpoBondDims(NUM_SITES - 1, MPO_BOND_DIM);

  cudensitymatMatrixProductOperator_t mpoHandle;
  std::vector<cudensitymatWrappedTensorCallback_t> mpoCallbacks(NUM_SITES, cudensitymatTensorCallbackNone);
  std::vector<cudensitymatWrappedTensorGradientCallback_t> mpoGradCallbacks(NUM_SITES, cudensitymatTensorGradientCallbackNone);
  HANDLE_CUDM_ERROR(cudensitymatCreateMatrixProductOperator(handle,
                      NUM_SITES,
                      spaceShape.data(),
                      CUDENSITYMAT_BOUNDARY_CONDITION_OPEN,
                      mpoBondDims.data(),
                      DATA_TYPE,
```

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

### bfi Example

Source: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html?highlight=proxy

Example usage of the bfi instruction: bfi.b32 d,a,b,start,len;

```ptx
bfi.b32  d,a,b,start,len;
```

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

### Full example with plan cache and tag

Source: https://docs.nvidia.com/cuda/cutensor/2.1.0/plan_cache.html

Complete C++ example that sets up a contraction, uses the plan cache, and sets a tag for a second loop. Includes error handling, GPU timing, and cleanup.

```cpp
  1 #include <stdlib.h>
  2 #include <stdio.h>
  3
  4 #include <unordered_map>
  5 #include <vector>
  6 #include <cassert>
  7
  8 #include <cuda_runtime.h>
  9 #include <cutensor.h>
 10
 11 #define HANDLE_ERROR(x)                                               \
 12 { const auto err = x;                                                 \
 13   if( err != CUTENSOR_STATUS_SUCCESS )                                \
 14   { printf("Error: %s\n", cutensorGetErrorString(err)); exit(-1); } \
 15 };
 16
 17 #define HANDLE_CUDA_ERROR(x)                                      \
 18 { const auto err = x;                                             \
 19   if( err != cudaSuccess )                                        \
 20   { printf("Error: %s\n", cudaGetErrorString(err)); exit(-1); } \
 21 };
 22
 23 struct GPUTimer
 24 {
 25     GPUTimer()
 26     {
 27         cudaEventCreate(&start_);
 28         cudaEventCreate(&stop_);
 29         cudaEventRecord(start_, 0);
 30     }
 31
 32     ~GPUTimer()
 33     {
 34         cudaEventDestroy(start_);
 35         cudaEventDestroy(stop_);
 36     }
 37
 38     void start()
 39     {
 40         cudaEventRecord(start_, 0);
 41     }
 42
 43     float seconds()
 44     {
 45         cudaEventRecord(stop_, 0);
 46         cudaEventSynchronize(stop_);
 47         float time;
 48         cudaEventElapsedTime(&time, start_, stop_);
 49         return time * 1e-3;
 50     }
 51     private:
 52     cudaEvent_t start_, stop_;
 53 };
 54
 55 int main()
 56 {
 57     typedef float floatTypeA;
 58     typedef float floatTypeB;
 59     typedef float floatTypeC;
 60     typedef float floatTypeCompute;
 61
 62     cutensorDataType_t typeA = CUTENSOR_R_32F;
 63     cutensorDataType_t typeB = CUTENSOR_R_32F;
 64     cutensorDataType_t typeC = CUTENSOR_R_32F;
 65     const cutensorComputeDescriptor_t descCompute = CUTENSOR_COMPUTE_DESC_32F;
 66
 67     floatTypeCompute alpha = (floatTypeCompute)1.1f;
 68     floatTypeCompute beta  = (floatTypeCompute)0.f;
 69
 70     /**********************
 71      * Computing: C_{m,u,n,v} = alpha * A_{m,h,k,n} B_{u,k,v,h} + beta * C_{m,u,n,v}
 72      **********************/
 73
 74     std::vector<int> modeC{'m','u','n','v'};
 75     std::vector<int> modeA{'m','h','k','n'};
 76     std::vector<int> modeB{'u','k','v','h'};
 77     int nmodeA = modeA.size();
 78     int nmodeB = modeB.size();
 79     int nmodeC = modeC.size();
 80
 81     std::unordered_map<int, int64_t> extent;
 82     extent['m'] = 96;
 83     extent['n'] = 96;
 84     extent['u'] = 96;
 85     extent['v'] = 64;
 86     extent['h'] = 64;
 87     extent['k'] = 64;
 88
 89     double gflops = (2.0 * extent['m'] * extent['n'] * extent['u'] * extent['v'] * extent['k'] * extent['h']) /1e9;
 90
 91     std::vector<int64_t> extentC;
 92     for (auto mode : modeC)
 93         extentC.push_back(extent[mode]);
 94     std::vector<int64_t> extentA;
 95     for (auto mode : modeA)
 96         extentA.push_back(extent[mode]);
 97     std::vector<int64_t> extentB;
 98     for (auto mode : modeB)
 99         extentB.push_back(extent[mode]);
100
101     /**********************
102      * Allocating data
103      **********************/
104
105     size_t elementsA = 1;
106     for (auto mode : modeA)
107         elementsA *= extent[mode];
108     size_t elementsB = 1;
109     for (auto mode : modeB)
110         elementsB *= extent[mode];
111     size_t elementsC = 1;
112     for (auto mode : modeC)
113         elementsC *= extent[mode];
114
```

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

### Run G5: First pipeline

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

Command to run the introduction pipeline example.

```bash
./build/example/cublasdx/01_gemm_introduction/introduction_pipeline
```

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

### Install CUDA Toolkit via RPM Installer

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

Install the repository meta-data, clean the yum cache, and install the CUDA toolkit. The driver must be installed separately starting with CUDA 13.4.

```bash
sudo rpm --install cuda-repo-<distro>-<version>.<architecture>.rpm
sudo rpm --erase gpg-pubkey-7fa2af80*
sudo yum clean expire-cache
sudo yum install cuda-toolkit
```
Tokens:
Raw