NVIDIA CUDA cuBLAS

https://docs.nvidia.com/cuda/cublas/
Admin
NVIDIA CUDA cuBLAS is a GPU-accelerated library providing optimized BLAS routines for linear algebra...
Tokens:448,495
Snippets:1,608
Trust Score:10
Update:2 weeks ago
Show doc for...
Tokens:
Raw
### Run cuBLASDx example executables

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, or timings and throughput for performance.

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

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

### Build a single cuBLASDx example target

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

Use the CMake target name to build one example. The target keeps the numeric prefix and drops the duplicate part when the directory and file name repeat words.

```bash
cmake --build build --target 11_gemm_device_performance
```

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

### Run introduction_pipeline

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

Run the introduction_pipeline executable to see host pipeline setup and device handle use.

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

```

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

### Run device_gemm_performance

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

Run the device_gemm_performance executable for a full-device GEMM tiling example.

```bash
./build/example/cublasdx/11_gemm_device_performance/device_gemm_performance

```

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

### Run simple_gemm_fp32

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

Run the simple_gemm_fp32 executable for a first checked fp32 GEMM example.

```bash
./build/example/cublasdx/02_gemm_precisions/simple_gemm_fp32

```

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

### Run tensor views TRSM example

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

Run command for the tensor views TRSM example. Executes the compiled binary that demonstrates applying conj_transpose_view at execute time so TRSM can solve against a transposed or conjugate-transposed triangular matrix without changing storage.

```bash
./build/example/cublasdx/18_tensor_transform/trsm_conj_transpose
```

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

### Get supported Tile IDs with cublasLtMatmulAlgoCapGetAttribute()

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

Example usage to retrieve the list of supported Tile IDs. The function writes the tile IDs into the provided buffer and returns the number of tiles via size_written.

```c
cublasLtMatmulTile_t tiles[CUBLASLT_MATMUL_TILE_END];
size_t num_tiles, size_written;
if (cublasLtMatmulAlgoCapGetAttribute(algo, CUBLASLT_ALGO_CAP_TILE_IDS, tiles, sizeof(tiles), &size_written) == CUBLAS_STATUS_SUCCESS) {
  num_tiles = size_written / sizeof(tiles[0]);}
```

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

### Run introduction_example

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

Run the introduction_example executable to see descriptor composition and launch basics for cuBLASDx GEMM.

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

```

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

### Configure and build cuBLASDx examples with CMake

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

Configure the build with BUILD_EXAMPLE=ON and set CUBLASDX_CUDA_ARCHITECTURES to match your GPU (default is SM80). Then build the aggregate target cublasdx_examples.

```bash
cmake -S . -B build -DBUILD_EXAMPLE=ON -DCUBLASDX_CUDA_ARCHITECTURES=80-real
cmake --build build --target cublasdx_examples -j
```

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

### Complete accumulator register GEMM kernel with tensor creation and copying

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

Complete example of the accumulator register GEMM kernel after filling in tensor creation and copying steps. It creates global and shared memory tensors, copies data from global to shared memory, waits for the copy, gets the default accumulator, executes the GEMM, and stores results back to global memory using partition_and_copy.

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

template<class GEMM>
__global__ void gemm_kernel_registers_accumulation(const typename GEMM::a_value_type* a,
                                                   const typename GEMM::b_value_type* b,
                                                   typename GEMM::c_value_type* c) {
    extern __shared__ __align__(16) cublasdx::byte smem[];

    // Make global memory tensor
    auto a_global_tensor = cublasdx::make_tensor(a, GEMM::get_layout_gmem_a());
    auto b_global_tensor = cublasdx::make_tensor(b, GEMM::get_layout_gmem_b());
    auto c_global_tensor = cublasdx::make_tensor(c, GEMM::get_layout_gmem_c());

    // Make shared memory tensor
    auto [smem_a, smem_b] = cublasdx::slice_shared_memory_ab<GEMM>(smem);
    auto a_shared_tensor = cublasdx::make_tensor(smem_a, GEMM::get_layout_smem_a());
    auto b_shared_tensor = cublasdx::make_tensor(smem_b, GEMM::get_layout_smem_b());

    // Load data from global memory tensor to shared memory tensor
    using alignment = cublasdx::alignment_of<GEMM>;
    cublasdx::copy<GEMM, alignment::a>(a_global_tensor, a_shared_tensor);
    cublasdx::copy<GEMM, alignment::b>(b_global_tensor, b_shared_tensor);
    cublasdx::copy_wait();

    // Get default data accumulator
    auto accumulator = GEMM::get_accumulator();

    // Execute GEMM with accumulation
    GEMM().execute(a_shared_tensor, b_shared_tensor, accumulator);

    // Partition Global C for GEMM and store appropriate elements to global memory
    auto results = accumulator.get_results();
    // Perform any epilogue on result data
    // ...

    // Store data back
    accumulator.partition_and_copy(results, c_global_tensor);
}
```

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

### TRSM Thread Execute Example

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

Example of using the TRSM thread execute method with a 32x16 matrix, float precision, real type, left side, lower fill mode, non-unit diagonal, SM 900, and Thread operator. It retrieves per-thread batches from global memory using get_batch and calls execute on the TRSM object. Refer to example 17 in the Examples section for a complete 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::Thread());

auto batch_a = cublasdx::get_batch(global_a, TRSM::get_layout_gmem_a(), threadIdx.x);
auto batch_b = cublasdx::get_batch(global_b, TRSM::get_layout_gmem_b(), threadIdx.x);
TRSM{}.execute(batch_a, batch_b);
```

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

### introduction_example.cu source

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

Source file for the introduction_example, showing GEMM descriptor composition from Size, Precision, Arrangement, Function, SM, and Block, and the three GEMM execution styles.

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

```

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

### Dynamic shared memory setup for large GEMM/TRSM operations

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

Defines a GEMM type with fp8 inputs and float output, then shows how to query required shared memory sizes using cublasdx::get_shared_storage_size (for GEMM) and cublasdx::get_shared_storage_size_ab (for TRSM and register API GEMM). It also demonstrates setting the max dynamic shared memory attribute with cudaFuncSetAttribute and launching the kernel with the required shared memory size. Note that the example includes a placeholder '(...)' for omitted code.

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

using GEMM = decltype(cublasdx::Size<128, 128, 64>()
              + cublasdx::Precision<__nv_fp8_e4m3, __nv_fp8_e5m2, float>()
              + cublasdx::Type<cublasdx::type::real>()
              + cublasdx::Arrangement<cublasdx::row_major, cublasdx::col_major>()
              + cublasdx::Function</* MM or TRSM */>()
              + cublasdx::SM<900>()
              + cublasdx::Block()); // Thread() will output single problem requirement for TRSM

void example() {
  (...)

  // Get required shared memory sizes, options:

  // Shared Memory API GEMM
  // 1 - Shared memory size required for matrices based on GEMM definition
  auto shared_memory_size = cublasdx::get_shared_storage_size<GEMM>();
  // 2 - Shared memory size when dynamic leading dimensions are used
  auto shared_memory_size = cublasdx::get_shared_storage_size<GEMM>(lda, ldb, ldc);
  // 3 - Shared memory size calculated based on custom matrix layouts for A, B, C matrices
  auto shared_memory_size = cublasdx::get_shared_storage_size<GEMM>(matrix_a_layout, matrix_b_layout, matrix_c_layout);

  // TRSM and register API GEMM
  // 1 - Shared memory size required for matrices based on GEMM definition
  auto shared_memory_size = cublasdx::get_shared_storage_size_ab<GEMM>();
  // 2 - Shared memory size when dynamic leading dimensions are used
  auto shared_memory_size = cublasdx::get_shared_storage_size_ab<GEMM>(lda, ldb);
  // 3 - Shared memory size calculated based on custom matrix layouts for A, B matrices
  auto shared_memory_size = cublasdx::get_shared_storage_size_ab<GEMM>(matrix_a_layout, matrix_b_layout);

  // Increases the max dynamic shared memory size to match GEMM requirements
  cudaFuncSetAttribute(gemm_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, shared_memory_size)
  // Invokes kernel with GEMM::block_dim threads in CUDA block
  gemm_kernel<GEMM><<<1, GEMM::block_dim, shared_memory_size>>>(alpha, a, b, beta, c);

  (...)
}
```

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

### introduction_pipeline.cu source

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

Source file for introduction_pipeline, showing host setup: global tensor descriptors, pipeline creation, device handle extraction, dynamic shared-memory sizing, asynchronous load and compute stage overlap, and final epilogue.

```cuda
example/cublasdx/01_gemm_introduction/introduction_pipeline.cu

```

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

### Example: Using suggest_layout_smem_a/b/c() with make_tensor

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

Shows how to define a BLAS type, slice shared memory, and create tensors with suggested layouts for optimal performance. The suggested layouts ignore leading dimensions and are best for row-major A and column-major B.

```cpp
using BLAS = decltype(Size<128, 128, 128>() + Type<type::real>() + Block() + Precision<float, float, double>());

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

// Slice shared memory into pointer for A, B, and C matrices
auto [smem_a, smem_b, smem_c] = cublasdx::slice_shared_memory<BLAS>(smem);

// Create suggested shared memory layout for optimal performance
auto suggested_smem_a = cublasdx::make_tensor(smem_a, BLAS::suggest_layout_smem_a());
auto suggested_smem_b = cublasdx::make_tensor(smem_b, BLAS::suggest_layout_smem_b());
auto suggested_smem_c = cublasdx::make_tensor(smem_c, BLAS::suggest_layout_smem_c());
```

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

### is_supported_rmem_restrict example

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

Example showing how to use is_supported_rmem_restrict with a BLAS description that has size, function, transpose mode, and Block operator but no SM operator. For the given BLAS and architecture 800, is_supported_smem_restrict returns false while is_supported_rmem_restrict returns true.

```cpp
using namespace cublasdx;

using BLAS = decltype(Size<128, 128, 128>() + Type<type::real>() + Block() + Precision<float>());
cublasdx::is_supported_smem_restrict<BLAS, 800>::value; // false
cublasdx::is_supported_rmem_restrict<BLAS, 800>::value; // true
```

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

### Run thread-level TRSM example

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

Run command for the thread-level TRSM example. Executes the compiled binary that demonstrates each CUDA thread solving one independent triangular system directly from global memory using the Thread() execution operator.

```bash
./build/example/cublasdx/17_trsm/trsm_thread
```

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

### Accumulator example #1

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

Demonstrates how to obtain the accumulator, create an empty register fragment, and copy global memory elements into it before executing a GEMM operation.

```APIDOC
## [C++] Accumulator example #1

### Description
Demonstrates how to obtain the accumulator, create an empty register fragment, and copy global memory elements into it before executing a GEMM operation.

### Method
C++ code snippet

### Endpoint
N/A (C++ code snippet)

### Parameters
N/A

### Request Example
```cpp
auto c_global_tensor = ...;

// Get accumulator
auto accumulator = BLAS::get_accumulator();
// Create local register fragment
auto d_register_fragment = accumulator.make_empty_fragment();
// Get view of this thread's global memory subtensor
auto c_global_partition = accumulator.partition_like_C(c_global_tensor);

// Ensure that all elements are in-bounds and
// no predication is necessary
static_assert(not accumulator.is_predicated());

// If this thread takes part in GEMM
if(accumulator.is_thread_active()) {
   // For each element of register fragment
   for(int i = 0; i < cublasdx::size(d_register_fragment); ++i) {
      // Copy respective global element into it
      d_register_fragment(i) = c_global_partition(i);
   }
}

// C += A * B
BLAS().execute(a_shared_tensor, b_shared_tensor, accumulator);

cublasdx::axpby(alpha, accumulator.get_results, beta, d_register_fragment);
```

### Response
#### Success Response
No explicit response is defined; the code snippet demonstrates usage of the accumulator API.

#### Response Example
N/A
```

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

### cublasXtCreate()

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

Initializes the cuBLASXt API and creates a handle to an opaque structure holding the cuBLASXt API context. It allocates hardware resources on the host and device and must be called prior to making any other cuBLASXt API calls.

```APIDOC
## cublasXtCreate()

### Description
Initializes the cuBLASXt API and creates a handle to an opaque structure holding the cuBLASXt API context. It allocates hardware resources on the host and device and must be called prior to making any other cuBLASXt API calls.

### Signature
```c
cublasStatus_t
cublasXtCreate(cublasXtHandle_t *handle)
```

### Parameters
- **handle** (cublasXtHandle_t*) - Output - Pointer to the handle to be created.

### Return Values
- `CUBLAS_STATUS_SUCCESS` - the initialization succeeded
- `CUBLAS_STATUS_ALLOC_FAILED` - the resources could not be allocated
- `CUBLAS_STATUS_NOT_SUPPORTED` - cuBLASXt API is only supported on 64-bit platform
```
Tokens:
Raw