default_gemv.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2017-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 "cutlass/gemm/threadblock/gemv.h"
#include "cutlass/gemm/threadblock/default_gemv_core.h"
#include "cutlass/gemm/threadblock/threadblock_swizzle.h"
namespace cutlass {
namespace gemm {
namespace kernel {
template <
typename ThreadBlockShape_,
typename ThreadShape_,
typename ElementA_,
typename LayoutA_,
typename ElementB_,
typename LayoutB_,
typename ElementCD_,
typename LayoutCD_,
typename ElementAccumulator_ = ElementCD_>
struct DefaultGemv {
using ThreadBlockShape = ThreadBlockShape_;
using ThreadShape = ThreadShape_;
using ElementA = ElementA_;
using LayoutA = LayoutA_;
using ElementB = ElementB_;
using LayoutB = LayoutB_;
using ElementAccumulator = ElementAccumulator_;
using LayoutAccumulator = LayoutCD_;
using ElementCD = ElementCD_;
using LayoutCD = LayoutCD_;
// Define the core components
using Core = typename cutlass::gemm::threadblock::DefaultGemvCore<
ThreadBlockShape, ThreadShape, ElementA, LayoutA, ElementB, LayoutB,
ElementAccumulator, LayoutAccumulator>;
// Define the threadblock-scoped gemv
using ThreadBlockGemv = cutlass::gemm::threadblock::Gemv<Core>;
// Iterator for multiplicand A
using IteratorA = typename ThreadBlockGemv::IteratorA;
// Iterator for multiplicand B
using IteratorB = typename ThreadBlockGemv::IteratorB;
using IteratorPolicyCD = typename platform::conditional<
platform::is_same<LayoutCD, layout::RowMajor>::value,
cutlass::transform::PitchLinearTilePolicyStripminedThreadContiguous<
layout::PitchLinearShape<ThreadBlockShape::kN, ThreadBlockShape::kM>, Core::kThreadsPerN, ThreadShape::kN>,
cutlass::transform::PitchLinearTilePolicyStripminedThreadStrided<
layout::PitchLinearShape<ThreadBlockShape::kM, ThreadBlockShape::kN>, Core::kThreadsPerN, ThreadShape::kM>>::type;
using IteratorCD = cutlass::transform::threadblock::PredicatedTileIterator<
cutlass::MatrixShape<ThreadBlockShape::kM, ThreadBlockShape::kN>, ElementCD, LayoutCD, 0, IteratorPolicyCD>;
using FragmentCD = typename IteratorCD::Fragment;
// Define the threadblock swizzle
using ThreadBlockSwizzle = cutlass::gemm::threadblock::GemvBatchedStridedThreadblockDefaultSwizzle;
};
} // namespace kernel
} // namespace gemm
} // namespace cutlass
cutlass::MatrixShape
Describes the size of a matrix tile.
Definition: matrix_shape.h:42
cutlass
Definition: aligned_buffer.h:35
cutlass::platform::is_same
std::is_same (false specialization)
Definition: platform.h:394
cutlass::gemm::kernel::DefaultGemv::ElementAccumulator
ElementAccumulator_ ElementAccumulator
Data type of accumulators.
Definition: default_gemv.h:78
cutlass::gemm::kernel::DefaultGemv::IteratorB
typename ThreadBlockGemv::IteratorB IteratorB
Definition: default_gemv.h:101
cutlass::gemm::kernel::DefaultGemv::ThreadShape
ThreadShape_ ThreadShape
Shape of warp-level matrix operation (concept: GemmShape)
Definition: default_gemv.h:63
cutlass::gemm::kernel::DefaultGemv::LayoutCD
LayoutCD_ LayoutCD
Layout of input/output matrix C/D.
Definition: default_gemv.h:87
cutlass::gemm::kernel::DefaultGemv::LayoutAccumulator
LayoutCD_ LayoutAccumulator
Data type of accumulators (same as C/D)
Definition: default_gemv.h:81
cutlass::layout::PitchLinearShape
Template defining a shape used by pitch-linear operators.
Definition: pitch_linear.h:43
cutlass::gemm::kernel::DefaultGemv::ElementA
ElementA_ ElementA
Data type of multiplicand A.
Definition: default_gemv.h:66
cutlass::gemm::kernel::DefaultGemv::FragmentCD
typename IteratorCD::Fragment FragmentCD
Fragment storage for C/D.
Definition: default_gemv.h:116
cutlass::gemm::threadblock::Gemv
Structure to compute the matrix-vector product using SIMT math instructions.
Definition: gemv.h:50
cutlass::gemm::kernel::DefaultGemv::LayoutA
LayoutA_ LayoutA
Layout of multiplicand A.
Definition: default_gemv.h:69
cutlass::gemm::threadblock::DefaultGemvCore
Definition: default_gemv_core.h:68
cutlass::platform::conditional
std::conditional (true specialization)
Definition: platform.h:325
cutlass::gemm::kernel::DefaultGemv
Definition: default_gemv.h:57
cutlass::gemm::kernel::DefaultGemv::IteratorPolicyCD
typename platform::conditional< platform::is_same< LayoutCD, layout::RowMajor >::value, cutlass::transform::PitchLinearTilePolicyStripminedThreadContiguous< layout::PitchLinearShape< ThreadBlockShape::kN, ThreadBlockShape::kM >, Core::kThreadsPerN, ThreadShape::kN >, cutlass::transform::PitchLinearTilePolicyStripminedThreadStrided< layout::PitchLinearShape< ThreadBlockShape::kM, ThreadBlockShape::kN >, Core::kThreadsPerN, ThreadShape::kM >>::type IteratorPolicyCD
Policy for the iterator that reads/writes C/D.
Definition: default_gemv.h:109
cutlass::gemm::threadblock::Gemv::IteratorA
typename Core_::IteratorA IteratorA
Iterates over A in global memory.
Definition: gemv.h:58
cutlass::transform::threadblock::PredicatedTileIterator
Definition: transform/threadblock/predicated_tile_iterator.h:133
default_gemv_core.h
Defines basic properties needed by CTA-level batched GEMV assuming expectations about data layout of ...
cutlass::gemm::threadblock::Gemv::IteratorB
typename Core_::IteratorB IteratorB
Iterates over B in global memory.
Definition: gemv.h:61
cutlass::gemm::threadblock::GemvBatchedStridedThreadblockDefaultSwizzle
Threadblock swizzling function for batched GEMVs.
Definition: gemm/threadblock/threadblock_swizzle.h:296
gemv.h
Template for a threadblock-scoped GEMV kernel.
cutlass::gemm::kernel::DefaultGemv::IteratorA
typename ThreadBlockGemv::IteratorA IteratorA
Definition: default_gemv.h:98
threadblock_swizzle.h
Implements several possible threadblock-swizzling functions mapping blockIdx to GEMM problems...
cutlass::gemm::kernel::DefaultGemv::Core
typename cutlass::gemm::threadblock::DefaultGemvCore< ThreadBlockShape, ThreadShape, ElementA, LayoutA, ElementB, LayoutB, ElementAccumulator, LayoutAccumulator > Core
Definition: default_gemv.h:92
cutlass::gemm::kernel::DefaultGemv::ElementB
ElementB_ ElementB
Data type of multiplicand B.
Definition: default_gemv.h:72
cutlass::transform::PitchLinearTilePolicyStripminedThreadStrided
Definition: pitch_linear_thread_map.h:168
cutlass::gemm::kernel::DefaultGemv::LayoutB
LayoutB_ LayoutB
Layout of multiplicand B.
Definition: default_gemv.h:75
cutlass::transform::PitchLinearTilePolicyStripminedThreadContiguous
Definition: pitch_linear_thread_map.h:140
cutlass::gemm::kernel::DefaultGemv::ThreadBlockShape
ThreadBlockShape_ ThreadBlockShape
Shape of Threadblock-level matrix operation (concept: GemmShape)
Definition: default_gemv.h:60
cutlass::gemm::kernel::DefaultGemv::ElementCD
ElementCD_ ElementCD
Data type of input/output matrix C/D.
Definition: default_gemv.h:84