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"
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_>
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_> {
64  typedef typename Base::Threads Threads;
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>
76  struct ThreadOffset {
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
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
Definition: convert.h:33
Defines iterators for efficiently loading and storing to global memory.
Definition: gemm_global_tile.h:70
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:368
Shape< 1, 2, Base::VectorizedTile::kC > ThreadsDelta
The threads strides.
Definition: hgemm_global_tile.h:66
Kind
Enumeration defining fundamental contiguous layouts.
Definition: matrix_traits.h:159
Base::Threads Threads
The threads.
Definition: hgemm_global_tile.h:64
static int const kH
The height of the cube.
Definition: shape.h:68
CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
Definition: hgemm_global_tile.h:78
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
Defines a type for restructuring a tile.
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ > Base
The base class.
Definition: hgemm_global_tile.h:62
Definition: hgemm_global_tile.h:48
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
The threads shape.
Definition: gemm_global_tile.h:88
static int const kW
The width of the cube.
Definition: shape.h:70
static int const kAccessSize
The number of scalars per LDG/STG.
Definition: gemm_global_tile.h:80
Computes the thread offset in (H, W) based on thread ID.
Definition: hgemm_global_tile.h:76
Kind
Definition: matrix_traits.h:357
Defines properties of matrices used to denote layout and operands to GEMM kernels.