Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

hgemm_global_tile.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Redistribution and use in source and binary forms, with or without modification, are permitted
5 * provided that the following conditions are met:
6 * * Redistributions of source code must retain the above copyright notice, this list of
7 * conditions and the following disclaimer.
8 * * Redistributions in binary form must reproduce the above copyright notice, this list of
9 * conditions and the following disclaimer in the documentation and/or other materials
10 * provided with the distribution.
11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12 * to endorse or promote products derived from this software without specific prior written
13 * permission.
14 *
15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23 *
24 **************************************************************************************************/
30 #pragma once
31
32 #include "cutlass/coord.h"
33 #include "cutlass/gemm/gemm_global_tile.h"
34 #include "cutlass/matrix_traits.h"
35 #include "cutlass/reshape_tile.h"
36
37 namespace cutlass {
38 namespace gemm {
39
41
42 template <GemmOperand::Kind kOperand_,
43 MatrixLayout::Kind kLayout_,
44 typename Scalar_,
45 typename Tile_,
46 typename Threads_,
47 int kAccessSize_>
48 struct HgemmCrosswiseGlobalTileTraits : public GemmGlobalTileTraits<
49 // Which GEMM operand?
50 kOperand_,
51 // The layout.
52 kLayout_,
53 // The scalar.
54 Scalar_,
55 // The tile.
56 Tile_,
57 // The threads.
58 Threads_,
59 // The number of scalars per LDG/STG.
60 kAccessSize_> {
62 typedef GemmGlobalTileTraits<kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_> Base;
64 typedef typename Base::Threads Threads;
66 typedef Shape<1, 2, Base::VectorizedTile::kC> ThreadsDelta;
68 typedef Shape<Base::Threads::kH * 2, 1, Base::Threads::kW, Base::kAccessSize> Delta;
70 typedef Shape<Base::VectorizedTile::kH / Base::Threads::kH / 2,
71 2,
72 Base::VectorizedTile::kW / Base::Threads::kW,
73 Base::VectorizedTile::kC / Base::kAccessSize>
74 Iterations;
76 struct ThreadOffset {
77 CUTLASS_HOST_DEVICE
78 Coord<4> operator()() const {
79 int thread_offset_h = threadIdx.x / Threads::kW * ThreadsDelta::kH;
80 int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
81
82 return make_Coord(0, thread_offset_h, thread_offset_w, 0);
83 }
84 };
85 };
86
88
89 } // namespace gemm
90 } // namespace cutlass
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Iterations
Shape< Base::VectorizedTile::kH/Base::Threads::kH/2, 2, Base::VectorizedTile::kW/Base::Threads::kW, Base::VectorizedTile::kC/Base::kAccessSize > Iterations
The number of iterations needed to load/store the tile.
Definition: hgemm_global_tile.h:74
cutlass
Definition: convert.h:33
gemm_global_tile.h
Defines iterators for efficiently loading and storing to global memory.
cutlass::gemm::GemmGlobalTileTraits
Definition: gemm_global_tile.h:70
coord.h
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:368
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadsDelta
Shape< 1, 2, Base::VectorizedTile::kC > ThreadsDelta
The threads strides.
Definition: hgemm_global_tile.h:66
cutlass::MatrixLayout::Kind
Kind
Enumeration defining fundamental contiguous layouts.
Definition: matrix_traits.h:159
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Threads
Base::Threads Threads
The threads.
Definition: hgemm_global_tile.h:64
cutlass::Shape::kH
static int const kH
The height of the cube.
Definition: shape.h:68
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadOffset::operator()
CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
Definition: hgemm_global_tile.h:78
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Delta
Shape< Base::Threads::kH *2, 1, Base::Threads::kW, Base::kAccessSize > Delta
The strides in each dimension between different loads/stores.
Definition: hgemm_global_tile.h:68
reshape_tile.h
Defines a type for restructuring a tile.
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Base
GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ > Base
The base class.
Definition: hgemm_global_tile.h:62
cutlass::gemm::HgemmCrosswiseGlobalTileTraits
Definition: hgemm_global_tile.h:48
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
cutlass::Coord< 4 >
cutlass::gemm::GemmGlobalTileTraits::Threads
ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
The threads shape.
Definition: gemm_global_tile.h:88
cutlass::Shape::kW
static int const kW
The width of the cube.
Definition: shape.h:70
cutlass::gemm::GemmGlobalTileTraits::kAccessSize
static int const kAccessSize
The number of scalars per LDG/STG.
Definition: gemm_global_tile.h:80
cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadOffset
Computes the thread offset in (H, W) based on thread ID.
Definition: hgemm_global_tile.h:76
cutlass::GemmOperand::Kind
Kind
Definition: matrix_traits.h:357
matrix_traits.h
Defines properties of matrices used to denote layout and operands to GEMM kernels.