library.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2019, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without modification, are permitted
* provided that the following conditions are met:
* * Redistributions of source code must retain the above copyright notice, this list of
* conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright notice, this list of
* conditions and the following disclaimer in the documentation and/or other materials
* provided with the distribution.
* * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
* to endorse or promote products derived from this software without specific prior written
* permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
* IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
* FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
* STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
**************************************************************************************************/
#pragma once
#include <vector>
#include <string>
#include <cstdint>
#include <cuda_runtime.h>
#include "cutlass/cutlass.h"
#include "cutlass/matrix_coord.h"
#include "cutlass/tensor_coord.h"
#include "cutlass/layout/tensor.h"
#include "cutlass/gemm/gemm.h"
namespace cutlass {
namespace library {
enum class LayoutTypeID {
kUnknown,
kColumnMajor,
kRowMajor,
kColumnMajorInterleavedK4,
kRowMajorInterleavedK4,
kColumnMajorInterleavedK16,
kRowMajorInterleavedK16,
kTensorNCHW,
kTensorNHWC,
kInvalid
};
enum class NumericTypeID {
kUnknown,
kVoid,
kB1,
kU4,
kU8,
kU16,
kU32,
kU64,
kS4,
kS8,
kS16,
kS32,
kS64,
kF16,
kF32,
kF64,
kCF16,
kCF32,
kCF64,
kCS4,
kCS8,
kCS16,
kCS32,
kCS64,
kCU4,
kCU8,
kCU16,
kCU32,
kCU64,
kInvalid
};
enum class ComplexTransform {
kNone,
kConjugate
};
enum class OperationKind {
kGemm,
kInvalid
};
enum class ScalarPointerMode {
kHost,
kDevice,
kInvalid
};
enum class SplitKMode {
kNone,
kSerial,
kParallel,
kParallelSerial,
kInvalid
};
enum class OpcodeClassID {
kSimt,
kTensorOp,
kWmmaTensorOp,
kInvalid
};
enum class GemmKind {
kGemm,
kBatched,
kArray,
kPlanarComplex,
kPlanarComplexBatched,
kInvalid
};
template <typename T> T from_string(std::string const &);
char const *to_string(OperationKind type, bool pretty = false);
template <> OperationKind from_string<OperationKind>(std::string const &str);
char const *to_string(NumericTypeID type, bool pretty = false);
template <> NumericTypeID from_string<NumericTypeID>(std::string const &str);
int sizeof_bits(NumericTypeID type);
bool is_complex_type(NumericTypeID type);
NumericTypeID get_real_type(NumericTypeID type);
bool is_integer_type(NumericTypeID type);
bool is_signed_type(NumericTypeID type);
bool is_signed_integer(NumericTypeID type);
bool is_unsigned_integer(NumericTypeID type);
bool is_float_type(NumericTypeID type);
char const *to_string(Status status, bool pretty = false);
char const *to_string(LayoutTypeID layout, bool pretty = false);
template <> LayoutTypeID from_string<LayoutTypeID>(std::string const &str);
int get_layout_stride_rank(LayoutTypeID layout_id);
char const *to_string(OpcodeClassID type, bool pretty = false);
template <>
OpcodeClassID from_string<OpcodeClassID>(std::string const &str);
std::string lexical_cast(int64_t int_value);
bool lexical_cast(std::vector<uint8_t> &bytes, NumericTypeID type, std::string const &str);
std::string lexical_cast(std::vector<uint8_t> &bytes, NumericTypeID type);
bool cast_from_int64(std::vector<uint8_t> &bytes, NumericTypeID type, int64_t src);
bool cast_from_uint64(std::vector<uint8_t> &bytes, NumericTypeID type, uint64_t src);
bool cast_from_double(std::vector<uint8_t> &bytes, NumericTypeID type, double src);
struct MathInstructionDescription {
cutlass::gemm::GemmCoord instruction_shape;
NumericTypeID element_accumulator;
OpcodeClassID opcode_class;
//
// Methods
//
MathInstructionDescription(
cutlass::gemm::GemmCoord instruction_shape = cutlass::gemm::GemmCoord(),
NumericTypeID element_accumulator = NumericTypeID::kInvalid,
OpcodeClassID opcode_class = OpcodeClassID::kInvalid
):
instruction_shape(instruction_shape), element_accumulator(element_accumulator), opcode_class(opcode_class) {}
};
struct TileDescription {
cutlass::gemm::GemmCoord threadblock_shape;
int threadblock_stages;
cutlass::gemm::GemmCoord warp_count;
MathInstructionDescription math_instruction;
int minimum_compute_capability;
int maximum_compute_capability;
//
// Methods
//
TileDescription(
cutlass::gemm::GemmCoord threadblock_shape = cutlass::gemm::GemmCoord(),
int threadblock_stages = 0,
cutlass::gemm::GemmCoord warp_count = cutlass::gemm::GemmCoord(),
MathInstructionDescription math_instruction = MathInstructionDescription(),
int minimum_compute_capability = 0,
int maximum_compute_capability = 0
):
threadblock_shape(threadblock_shape),
threadblock_stages(threadblock_stages),
warp_count(warp_count),
math_instruction(math_instruction),
minimum_compute_capability(minimum_compute_capability),
maximum_compute_capability(maximum_compute_capability) { }
};
struct OperationDescription {
char const * name;
OperationKind kind;
TileDescription tile_description;
//
// Methods
//
OperationDescription(
char const * name = "unknown",
OperationKind kind = OperationKind::kInvalid,
TileDescription const & tile_description = TileDescription()
):
name(name), kind(kind), tile_description(tile_description) { }
};
struct TensorDescription {
NumericTypeID element;
LayoutTypeID layout;
int alignment;
int log_extent_range;
int log_stride_range;
//
// Methods
//
TensorDescription(
NumericTypeID element = NumericTypeID::kInvalid,
LayoutTypeID layout = LayoutTypeID::kInvalid,
int alignment = 1,
int log_extent_range = 24,
int log_stride_range = 24
):
element(element),
layout(layout),
alignment(alignment),
log_extent_range(log_extent_range),
log_stride_range(log_stride_range) { }
};
struct GemmDescription : public OperationDescription {
GemmKind gemm_kind;
TensorDescription A;
TensorDescription B;
TensorDescription C;
NumericTypeID element_epilogue;
SplitKMode split_k_mode;
ComplexTransform transform_A;
ComplexTransform transform_B;
//
// Methods
//
GemmDescription(
GemmKind gemm_kind = GemmKind::kGemm,
TensorDescription const &A = TensorDescription(),
TensorDescription const &B = TensorDescription(),
TensorDescription const &C = TensorDescription(),
NumericTypeID element_epilogue = NumericTypeID::kInvalid,
SplitKMode split_k_mode = SplitKMode::kNone,
ComplexTransform transform_A = ComplexTransform::kNone,
ComplexTransform transform_B = ComplexTransform::kNone
):
gemm_kind(gemm_kind),
A(A),
B(B),
C(C),
element_epilogue(element_epilogue),
split_k_mode(split_k_mode),
transform_A(transform_A),
transform_B(transform_B) {}
};
class Operation {
public:
virtual ~Operation() { }
virtual OperationDescription const & description() const = 0;
virtual Status can_implement(
void const *configuration,
void const *arguments) const = 0;
virtual uint64_t get_host_workspace_size(
void const *configuration) const = 0;
virtual uint64_t get_device_workspace_size(
void const *configuration) const = 0;
virtual Status initialize(
void const *configuration,
void *host_workspace,
void *device_workspace,
cudaStream_t stream = nullptr) const = 0;
virtual Status run(
void const *arguments,
void *host_workspace,
void *device_workspace = nullptr,
cudaStream_t stream = nullptr) const = 0;
};
//
// OperationKind: Gemm
// GemmKind: Gemm
//
struct GemmConfiguration {
gemm::GemmCoord problem_size;
int64_t lda;
int64_t ldb;
int64_t ldc;
int64_t ldd;
int split_k_slices;
};
struct GemmArguments {
void const *A;
void const *B;
void const *C;
void *D;
void const *alpha;
void const *beta;
ScalarPointerMode pointer_mode;
};
//
// OperationKind: Gemm
// GemmKind: Batched
struct GemmBatchedConfiguration {
gemm::GemmCoord problem_size;
int64_t lda;
int64_t ldb;
int64_t ldc;
int64_t ldd;
int64_t batch_stride_A;
int64_t batch_stride_B;
int64_t batch_stride_C;
int64_t batch_stride_D;
int batch_count;
};
using GemmBatchedArguments = GemmArguments;
//
// OperationKind: Gemm
// GemmKind: Array
struct GemmArrayConfiguration {
gemm::GemmCoord problem_size;
int64_t const *lda;
int64_t const *ldb;
int64_t const *ldc;
int64_t const *ldd;
int batch_count;
};
struct GemmArrayArguments {
void const * const *A;
void const * const *B;
void const * const *C;
void * const *D;
void const *alpha;
void const *beta;
ScalarPointerMode pointer_mode;
};
//
// OperationKind: Gemm
// GemmKind: Planar complex
struct GemmPlanarComplexConfiguration {
gemm::GemmCoord problem_size;
int64_t lda;
int64_t ldb;
int64_t ldc;
int64_t ldd;
int64_t imag_stride_A;
int64_t imag_stride_B;
int64_t imag_stride_C;
int64_t imag_stride_D;
};
using GemmPlanarComplexArgments = GemmArguments;
//
// OperationKind: Gemm
// GemmKind: Planar complex batched
//
struct GemmPlanarComplexBatchedConfiguration {
gemm::GemmCoord problem_size;
int64_t lda;
int64_t ldb;
int64_t ldc;
int64_t ldd;
int64_t imag_stride_A;
int64_t imag_stride_B;
int64_t imag_stride_C;
int64_t imag_stride_D;
int64_t batched_stride_A;
int64_t batched_stride_B;
int64_t batched_stride_C;
int64_t batched_stride_D;
};
using GemmPlanarComplexBatchedArguments = GemmArguments;
} // namespace library
} // namespace cutlass
cutlass::library::GemmPlanarComplexBatchedConfiguration::lda
int64_t lda
Definition: library.h:609
cutlass::library::NumericTypeID::kCS8
cutlass::library::TensorDescription::alignment
int alignment
Alignment restriction on pointers, strides, and extents.
Definition: library.h:336
cutlass::library::NumericTypeID::kCU32
cutlass::library::GemmArrayArguments::A
void const *const * A
Definition: library.h:565
cutlass::library::GemmKind::kPlanarComplexBatched
cutlass::library::NumericTypeID::kCS16
cutlass::library::OpcodeClassID::kWmmaTensorOp
cutlass::library::Operation::~Operation
virtual ~Operation()
Definition: library.h:421
cutlass::library::OperationDescription
High-level description of an operation.
Definition: library.h:304
cutlass::library::GemmKind::kPlanarComplex
cutlass::library::LayoutTypeID::kColumnMajor
cutlass::library::to_string
char const * to_string(OperationKind type, bool pretty=false)
Converts a NumericType enumerant to a string.
cutlass
Definition: aligned_buffer.h:35
cutlass::library::is_complex_type
bool is_complex_type(NumericTypeID type)
Returns true if the numeric type is a complex data type or false if real-valued.
cutlass::library::GemmArrayArguments::D
void *const * D
Definition: library.h:568
cutlass::library::TensorDescription::layout
LayoutTypeID layout
Enumerant identifying the layout function for the tensor.
Definition: library.h:333
cutlass::library::GemmDescription::gemm_kind
GemmKind gemm_kind
Indicates the kind of GEMM performed.
Definition: library.h:367
cutlass::library::GemmPlanarComplexConfiguration::ldc
int64_t ldc
Definition: library.h:587
cutlass::library::GemmArguments
Arguments for GEMM.
Definition: library.h:477
cutlass::library::GemmKind::kArray
cutlass::library::GemmArrayConfiguration::batch_count
int batch_count
Definition: library.h:560
cutlass::library::ComplexTransform
ComplexTransform
Enumeraed type describing a transformation on a complex value.
Definition: library.h:111
cutlass::library::LayoutTypeID::kColumnMajorInterleavedK16
cutlass::library::GemmArrayArguments::C
void const *const * C
Definition: library.h:567
cutlass::library::GemmPlanarComplexConfiguration::problem_size
gemm::GemmCoord problem_size
Definition: library.h:583
cutlass::library::GemmArrayConfiguration
Configuration for batched GEMM in which multiple matrix products are computed.
Definition: library.h:551
cutlass::library::is_signed_integer
bool is_signed_integer(NumericTypeID type)
Returns true if numeric type is a signed integer.
cutlass::library::GemmKind
GemmKind
Enumeration indicating what kind of GEMM operation to perform.
Definition: library.h:149
cutlass::library::NumericTypeID::kCF64
cutlass::library::NumericTypeID::kCU64
cutlass::library::get_real_type
NumericTypeID get_real_type(NumericTypeID type)
Returns the real-valued type underlying a type (only different from &#39;type&#39; if complex) ...
cutlass::library::from_string< OperationKind >
OperationKind from_string< OperationKind >(std::string const &str)
Parses a NumericType enumerant from a string.
cutlass::gemm::GemmCoord
Definition: include/cutlass/gemm/gemm.h:94
cutlass::library::get_layout_stride_rank
int get_layout_stride_rank(LayoutTypeID layout_id)
Returns the rank of a layout&#39;s stride base on the LayoutTypeID.
cutlass::library::GemmBatchedConfiguration::ldb
int64_t ldb
Leading dimension of B matrix.
Definition: library.h:517
cutlass::library::NumericTypeID::kCF32
cutlass::library::NumericTypeID::kCU8
cutlass::library::ComplexTransform::kNone
cutlass::library::GemmArrayConfiguration::ldc
int64_t const * ldc
Definition: library.h:557
cutlass::library::OpcodeClassID::kTensorOp
cutlass::library::NumericTypeID::kS64
cutlass::library::GemmPlanarComplexBatchedConfiguration::batched_stride_B
int64_t batched_stride_B
Definition: library.h:620
cutlass::library::GemmPlanarComplexConfiguration
Complex valued GEMM in which real and imaginary parts are separated by a stride.
Definition: library.h:581
cutlass::library::TensorDescription::log_stride_range
int log_stride_range
log2() of the maximum value each relevant stride may have
Definition: library.h:342
cutlass::library::ComplexTransform::kConjugate
cutlass::library::NumericTypeID::kB1
gemm.h
Defines common types used for all GEMM-like operators.
cutlass::library::GemmDescription::transform_A
ComplexTransform transform_A
Transformation on A operand.
Definition: library.h:385
cutlass::library::NumericTypeID::kS16
cutlass::library::SplitKMode::kSerial
cutlass::library::GemmPlanarComplexBatchedConfiguration::imag_stride_B
int64_t imag_stride_B
Definition: library.h:615
cutlass::library::GemmDescription::GemmDescription
GemmDescription(GemmKind gemm_kind=GemmKind::kGemm, TensorDescription const &A=TensorDescription(), TensorDescription const &B=TensorDescription(), TensorDescription const &C=TensorDescription(), NumericTypeID element_epilogue=NumericTypeID::kInvalid, SplitKMode split_k_mode=SplitKMode::kNone, ComplexTransform transform_A=ComplexTransform::kNone, ComplexTransform transform_B=ComplexTransform::kNone)
Definition: library.h:394
cutlass::library::sizeof_bits
int sizeof_bits(NumericTypeID type)
Returns the size of a data type in bits.
cutlass::library::Operation
Base class for all device-wide operations.
Definition: library.h:418
cutlass::library::GemmPlanarComplexConfiguration::imag_stride_A
int64_t imag_stride_A
Definition: library.h:590
cutlass::library::NumericTypeID::kCS64
cutlass::library::from_string< NumericTypeID >
NumericTypeID from_string< NumericTypeID >(std::string const &str)
Parses a NumericType enumerant from a string.
cutlass::library::LayoutTypeID
LayoutTypeID
Layout type identifier.
Definition: library.h:63
cutlass::library::OpcodeClassID
OpcodeClassID
Indicates the classificaition of the math instruction.
Definition: library.h:139
cutlass::library::GemmPlanarComplexBatchedConfiguration::ldc
int64_t ldc
Definition: library.h:611
cutlass::library::lexical_cast
std::string lexical_cast(int64_t int_value)
Lexical cast from int64_t to string.
cutlass::library::GemmArguments::pointer_mode
ScalarPointerMode pointer_mode
Enumerant indicating whether alpha/beta point to host or device memory.
Definition: library.h:498
cutlass::library::GemmArrayConfiguration::ldb
int64_t const * ldb
Definition: library.h:556
cutlass::library::GemmArrayConfiguration::problem_size
gemm::GemmCoord problem_size
Definition: library.h:553
cutlass::library::GemmPlanarComplexBatchedConfiguration::batched_stride_C
int64_t batched_stride_C
Definition: library.h:621
cutlass::library::OperationDescription::OperationDescription
OperationDescription(char const *name="unknown", OperationKind kind=OperationKind::kInvalid, TileDescription const &tile_description=TileDescription())
Definition: library.h:318
cutlass::library::TileDescription::maximum_compute_capability
int maximum_compute_capability
Minimum compute capability (e.g. 70, 75) of a device eligible to run the operation.
Definition: library.h:281
cutlass::library::NumericTypeID::kS8
cutlass::library::GemmConfiguration
Configuration for basic GEMM operations.
Definition: library.h:455
cutlass::library::GemmPlanarComplexBatchedConfiguration::imag_stride_D
int64_t imag_stride_D
Definition: library.h:617
cutlass::library::MathInstructionDescription
Definition: library.h:238
cutlass::library::GemmArguments::B
void const * B
Pointer to B matrix.
Definition: library.h:483
cutlass::library::GemmPlanarComplexBatchedConfiguration::imag_stride_A
int64_t imag_stride_A
Definition: library.h:614
cutlass::library::NumericTypeID::kU16
cutlass::library::LayoutTypeID::kRowMajorInterleavedK16
cutlass::library::GemmDescription::A
TensorDescription A
Describes the A operand.
Definition: library.h:370
cutlass::library::TileDescription
Structure describing the tiled structure of a GEMM-like computation.
Definition: library.h:263
cutlass::library::GemmConfiguration::split_k_slices
int split_k_slices
Number of partitions of K dimension.
Definition: library.h:473
cutlass::library::GemmPlanarComplexConfiguration::imag_stride_B
int64_t imag_stride_B
Definition: library.h:591
cutlass::library::from_string< OpcodeClassID >
OpcodeClassID from_string< OpcodeClassID >(std::string const &str)
Converts a OpcodeClassID enumerant from a string.
cutlass::library::NumericTypeID::kU8
cutlass::library::GemmArguments::A
void const * A
Pointer to A matrix.
Definition: library.h:480
tensor.h
Defines layout functions used by TensorRef and derived classes for common 4-D and 5-D tensor formats...
cutlass::library::GemmConfiguration::ldd
int64_t ldd
Leading dimension of D matrix.
Definition: library.h:470
cutlass::library::GemmDescription::transform_B
ComplexTransform transform_B
Transformation on B operand.
Definition: library.h:388
cutlass::library::GemmPlanarComplexBatchedConfiguration::ldb
int64_t ldb
Definition: library.h:610
cutlass::library::is_signed_type
bool is_signed_type(NumericTypeID type)
Returns true if numeric type is signed.
cutlass::library::NumericTypeID::kCU4
cutlass::library::GemmDescription::element_epilogue
NumericTypeID element_epilogue
Describes the data type of the scalars passed to the epilogue.
Definition: library.h:379
cutlass::library::GemmArrayConfiguration::lda
int64_t const * lda
Definition: library.h:555
cutlass::library::SplitKMode::kParallel
cutlass::library::GemmArrayConfiguration::ldd
int64_t const * ldd
Definition: library.h:558
cutlass::library::TileDescription::minimum_compute_capability
int minimum_compute_capability
Minimum compute capability (e.g. 70, 75) of a device eligible to run the operation.
Definition: library.h:278
cutlass::library::GemmPlanarComplexConfiguration::ldd
int64_t ldd
Definition: library.h:588
cutlass::library::GemmBatchedConfiguration::batch_stride_C
int64_t batch_stride_C
Stride between instances of the C matrix in memory.
Definition: library.h:532
cutlass::library::GemmArrayArguments::B
void const *const * B
Definition: library.h:566
cutlass::library::NumericTypeID
NumericTypeID
Numeric data type.
Definition: library.h:77
cutlass::library::GemmPlanarComplexConfiguration::lda
int64_t lda
Definition: library.h:585
cutlass::library::TileDescription::warp_count
cutlass::gemm::GemmCoord warp_count
Number of warps in each logical dimension.
Definition: library.h:272
cutlass::library::GemmConfiguration::lda
int64_t lda
Leading dimension of A matrix.
Definition: library.h:461
cutlass::library::is_float_type
bool is_float_type(NumericTypeID type)
Returns true if numeric type is floating-point type.
cutlass::library::LayoutTypeID::kColumnMajorInterleavedK4
cutlass::library::TileDescription::TileDescription
TileDescription(cutlass::gemm::GemmCoord threadblock_shape=cutlass::gemm::GemmCoord(), int threadblock_stages=0, cutlass::gemm::GemmCoord warp_count=cutlass::gemm::GemmCoord(), MathInstructionDescription math_instruction=MathInstructionDescription(), int minimum_compute_capability=0, int maximum_compute_capability=0)
Definition: library.h:287
cutlass::library::cast_from_double
bool cast_from_double(std::vector< uint8_t > &bytes, NumericTypeID type, double src)
Casts from a real value represented as a double to the destination type. Returns true if successful...
cutlass::library::NumericTypeID::kU32
cutlass::library::MathInstructionDescription::element_accumulator
NumericTypeID element_accumulator
Describes the data type of the internal accumulator.
Definition: library.h:244
tensor_coord.h
Defines a canonical coordinate for rank=4 tensors offering named indices.
cutlass::library::GemmDescription::B
TensorDescription B
Describes the B operand.
Definition: library.h:373
cutlass::library::LayoutTypeID::kInvalid
cutlass::library::GemmArrayArguments::alpha
void const * alpha
Definition: library.h:569
cutlass::library::NumericTypeID::kU4
cutlass::library::GemmArguments::beta
void const * beta
Host or device pointer to beta scalar.
Definition: library.h:495
cutlass::library::LayoutTypeID::kTensorNCHW
cutlass::library::GemmBatchedConfiguration::ldd
int64_t ldd
Leading dimension of D matrix.
Definition: library.h:523
cutlass::library::MathInstructionDescription::opcode_class
OpcodeClassID opcode_class
Classification of math instruction.
Definition: library.h:247
cutlass::library::NumericTypeID::kCF16
cutlass::library::GemmBatchedConfiguration::problem_size
gemm::GemmCoord problem_size
GEMM problem size.
Definition: library.h:511
cutlass::library::GemmConfiguration::ldc
int64_t ldc
Leading dimension of C matrix.
Definition: library.h:467
cutlass::library::LayoutTypeID::kTensorNHWC
cutlass::library::ScalarPointerMode::kHost
cutlass::library::GemmArguments::D
void * D
Pointer to D matrix.
Definition: library.h:489
cutlass::library::cast_from_uint64
bool cast_from_uint64(std::vector< uint8_t > &bytes, NumericTypeID type, uint64_t src)
Casts from an unsigned int64 to the destination type. Returns true if successful. ...
cutlass::library::GemmDescription::C
TensorDescription C
Describes the source and destination matrices.
Definition: library.h:376
cutlass::library::GemmPlanarComplexConfiguration::ldb
int64_t ldb
Definition: library.h:586
cutlass::library::GemmPlanarComplexBatchedConfiguration::imag_stride_C
int64_t imag_stride_C
Definition: library.h:616
cutlass::library::GemmPlanarComplexBatchedConfiguration
Batched complex valued GEMM in which real and imaginary parts are separated by a stride.
Definition: library.h:605
cutlass::library::OperationKind::kGemm
cutlass::library::GemmPlanarComplexBatchedConfiguration::batched_stride_D
int64_t batched_stride_D
Definition: library.h:622
cutlass::library::GemmBatchedConfiguration
Configuration for batched GEMM in which multiple matrix products are computed.
Definition: library.h:508
cutlass::library::GemmKind::kBatched
cutlass::library::LayoutTypeID::kRowMajor
cutlass::library::GemmBatchedConfiguration::batch_stride_A
int64_t batch_stride_A
Stride between instances of the A matrix in memory.
Definition: library.h:526
cutlass::library::is_integer_type
bool is_integer_type(NumericTypeID type)
Returns true if numeric type is integer.
cutlass::library::ScalarPointerMode
ScalarPointerMode
Enumeration indicating whether scalars are in host or device memory.
Definition: library.h:123
cutlass::library::TensorDescription::element
NumericTypeID element
Numeric type of an individual element.
Definition: library.h:330
cutlass::library::NumericTypeID::kCU16
cutlass::library::GemmBatchedConfiguration::batch_count
int batch_count
Number of GEMMs in batch.
Definition: library.h:538
cutlass::library::GemmArguments::C
void const * C
Pointer to C matrix.
Definition: library.h:486
cutlass::library::from_string
T from_string(std::string const &)
Lexical cast from string.
cutlass::library::TileDescription::threadblock_stages
int threadblock_stages
Describes the number of pipeline stages in the threadblock-scoped mainloop.
Definition: library.h:269
matrix_coord.h
Defines a canonical coordinate for rank=2 matrices offering named indices.
cutlass::library::GemmPlanarComplexConfiguration::imag_stride_D
int64_t imag_stride_D
Definition: library.h:593
cutlass::library::NumericTypeID::kS4
cutlass::library::NumericTypeID::kS32
cutlass::library::from_string< LayoutTypeID >
LayoutTypeID from_string< LayoutTypeID >(std::string const &str)
Parses a LayoutType enumerant from a string.
cutlass::library::GemmArrayArguments::pointer_mode
ScalarPointerMode pointer_mode
Definition: library.h:571
cutlass::library::MathInstructionDescription::MathInstructionDescription
MathInstructionDescription(cutlass::gemm::GemmCoord instruction_shape=cutlass::gemm::GemmCoord(), NumericTypeID element_accumulator=NumericTypeID::kInvalid, OpcodeClassID opcode_class=OpcodeClassID::kInvalid)
Definition: library.h:253
cutlass::library::GemmPlanarComplexBatchedConfiguration::batched_stride_A
int64_t batched_stride_A
Definition: library.h:619
cutlass::library::GemmDescription
Description of all GEMM computations.
Definition: library.h:364
cutlass::library::GemmBatchedConfiguration::lda
int64_t lda
Leading dimension of A matrix.
Definition: library.h:514
cutlass::library::GemmConfiguration::problem_size
gemm::GemmCoord problem_size
GEMM problem size.
Definition: library.h:458
cutlass::library::GemmDescription::split_k_mode
SplitKMode split_k_mode
Describes the structure of parallel reductions.
Definition: library.h:382
cutlass::library::cast_from_int64
bool cast_from_int64(std::vector< uint8_t > &bytes, NumericTypeID type, int64_t src)
Casts from a signed int64 to the destination type. Returns true if successful.
cutlass::library::TensorDescription::log_extent_range
int log_extent_range
log2() of the maximum extent of each dimension
Definition: library.h:339
cutlass::library::OperationDescription::name
char const * name
Unique identifier describing the operation.
Definition: library.h:307
cutlass::library::GemmBatchedConfiguration::batch_stride_B
int64_t batch_stride_B
Stride between instances of the B matrix in memory.
Definition: library.h:529
cutlass::library::MathInstructionDescription::instruction_shape
cutlass::gemm::GemmCoord instruction_shape
Shape of the target math instruction.
Definition: library.h:241
cutlass::library::GemmArrayArguments::beta
void const * beta
Definition: library.h:570
cutlass::library::OperationDescription::tile_description
TileDescription tile_description
Describes the tiled structure of a GEMM-like computation.
Definition: library.h:313
cutlass::library::NumericTypeID::kF16
cutlass::library::NumericTypeID::kF32
cutlass::library::GemmBatchedConfiguration::ldc
int64_t ldc
Leading dimension of C matrix.
Definition: library.h:520
cutlass::library::TensorDescription
Structure describing the properties of a tensor.
Definition: library.h:327
cutlass::library::GemmConfiguration::ldb
int64_t ldb
Leading dimension of B matrix.
Definition: library.h:464
cutlass::library::GemmPlanarComplexBatchedConfiguration::problem_size
gemm::GemmCoord problem_size
Definition: library.h:607
cutlass::library::is_unsigned_integer
bool is_unsigned_integer(NumericTypeID type)
returns true if numeric type is an unsigned integer
cutlass::library::GemmArrayArguments
Arguments for GEMM - used by all the GEMM operations.
Definition: library.h:564
cutlass::library::NumericTypeID::kCS32
cutlass::library::LayoutTypeID::kUnknown
cutlass::library::OpcodeClassID::kSimt
cutlass::library::OperationKind
OperationKind
Enumeration indicating the kind of operation.
Definition: library.h:117
cutlass::library::GemmArguments::alpha
void const * alpha
Host or device pointer to alpha scalar.
Definition: library.h:492
cutlass::library::SplitKMode::kParallelSerial
cutlass::library::OperationDescription::kind
OperationKind kind
Kind of operation.
Definition: library.h:310
cutlass::library::TileDescription::threadblock_shape
cutlass::gemm::GemmCoord threadblock_shape
Describes the shape of a threadblock (in elements)
Definition: library.h:266
cutlass::library::GemmPlanarComplexBatchedConfiguration::ldd
int64_t ldd
Definition: library.h:612
cutlass::library::NumericTypeID::kVoid
cutlass::library::ScalarPointerMode::kDevice
cutlass::library::NumericTypeID::kF64
cutlass::library::GemmPlanarComplexConfiguration::imag_stride_C
int64_t imag_stride_C
Definition: library.h:592
cutlass::library::TileDescription::math_instruction
MathInstructionDescription math_instruction
Core math instruction.
Definition: library.h:275
cutlass.h
Basic include for CUTLASS.
cutlass::library::SplitKMode
SplitKMode
Describes how reductions are performed across threadblocks.
Definition: library.h:130
cutlass::library::NumericTypeID::kU64
cutlass::library::NumericTypeID::kCS4
cutlass::Status
Status
Status code returned by CUTLASS operations.
Definition: cutlass.h:39
cutlass::library::GemmBatchedConfiguration::batch_stride_D
int64_t batch_stride_D
Stride between instances of the D matrix in memory.
Definition: library.h:535
cutlass::library::LayoutTypeID::kRowMajorInterleavedK4
cutlass::library::TensorDescription::TensorDescription
TensorDescription(NumericTypeID element=NumericTypeID::kInvalid, LayoutTypeID layout=LayoutTypeID::kInvalid, int alignment=1, int log_extent_range=24, int log_stride_range=24)
Definition: library.h:347