From 30f98c0674dfbedaf61dfe99a7afdd75069fc473 Mon Sep 17 00:00:00 2001 From: root Date: Sat, 15 Aug 2026 05:41:30 +0000 Subject: [PATCH] =?UTF-8?q?data:=20Cu10=20CUTLASS=20part=202=20=E2=80=94?= =?UTF-8?q?=20tensorop=20example,=20arch.h,=20cutlass.h?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- cat_files/CMakeLists_batched_gemm.txt | 27 + cat_files/CMakeLists_tensorop_gemm.txt | 27 + cat_files/arch.h | 84 +++ cat_files/arch_arch.h | 84 +++ cat_files/basic_gemm.cu | 492 +++++++++++++++++ cat_files/cutlass.h | 175 ++++++ cat_files/gemm_batched_full.h | 726 ++++++++++++++++++++++++ cat_files/gemm_device.h | 732 +++++++++++++++++++++++++ cat_files/turing_tensorop_gemm.cu | 354 ++++++++++++ 9 files changed, 2701 insertions(+) create mode 100644 cat_files/CMakeLists_batched_gemm.txt create mode 100644 cat_files/CMakeLists_tensorop_gemm.txt create mode 100644 cat_files/arch.h create mode 100644 cat_files/arch_arch.h create mode 100644 cat_files/basic_gemm.cu create mode 100644 cat_files/cutlass.h create mode 100644 cat_files/gemm_batched_full.h create mode 100644 cat_files/gemm_device.h create mode 100644 cat_files/turing_tensorop_gemm.cu diff --git a/cat_files/CMakeLists_batched_gemm.txt b/cat_files/CMakeLists_batched_gemm.txt new file mode 100644 index 00000000..6cd0ca8d --- /dev/null +++ b/cat_files/CMakeLists_batched_gemm.txt @@ -0,0 +1,27 @@ +# Copyright (c) 2017-2020, 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. + +cutlass_example_add_executable( + 05_batched_gemm + batched_gemm.cu +) + diff --git a/cat_files/CMakeLists_tensorop_gemm.txt b/cat_files/CMakeLists_tensorop_gemm.txt new file mode 100644 index 00000000..e920341a --- /dev/null +++ b/cat_files/CMakeLists_tensorop_gemm.txt @@ -0,0 +1,27 @@ +# Copyright (c) 2017-2020, 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. + +# cutlass_example_add_executable( +# 08_turing_tensorop_gemm +# turing_tensorop_gemm.cu +# ) + diff --git a/cat_files/arch.h b/cat_files/arch.h new file mode 100644 index 00000000..9bc09a23 --- /dev/null +++ b/cat_files/arch.h @@ -0,0 +1,84 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2021, 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. + * + **************************************************************************************************/ +/*! \file + \brief Defines tags for architecture-specific configurations. +*/ + +#pragma once + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +namespace cutlass { +namespace arch { + +//////////////////////////////////////////////////////////////////////////////////////////////////// +struct Sm50 { + static int const kMinComputeCapability = 50; +}; +struct Sm60 { + static int const kMinComputeCapability = 60; +}; +struct Sm61 { + static int const kMinComputeCapability = 61; +}; + + +/// BIGISLAND Arch +struct Cu10 { + static int const kMinComputeCapability = 10; +}; + +struct Sm62 { + static int const kMinComputeCapability = 62; +}; + +/// Triggers a breakpoint on the device +CUTLASS_DEVICE +void device_breakpoint() { +#if defined(__CUDA_ARCH__) + asm volatile (" brkpt;\n"); +#endif + +//////////////////////////////////////////////////////////////////////////////////////////////////// +/// Switches to control performance improvement only supported on Iluvatar platform + +/// Compiler of Iluvatar-CoreX implicitly convert boolean type that is stored at VRF to a 64-bit +/// width integer type on SRF +#define IMPLICIT_VRF_BOOLEAN_TO_SRF_INTEGER 1 + +/// Enable block load or store +#define BLOCK_LOAD_STORE 1 + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +} +//////////////////////////////////////////////////////////////////////////////////////////////////// + +} // namespace arch +} // namespace cutlass + +//////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/cat_files/arch_arch.h b/cat_files/arch_arch.h new file mode 100644 index 00000000..9bc09a23 --- /dev/null +++ b/cat_files/arch_arch.h @@ -0,0 +1,84 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2021, 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. + * + **************************************************************************************************/ +/*! \file + \brief Defines tags for architecture-specific configurations. +*/ + +#pragma once + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +namespace cutlass { +namespace arch { + +//////////////////////////////////////////////////////////////////////////////////////////////////// +struct Sm50 { + static int const kMinComputeCapability = 50; +}; +struct Sm60 { + static int const kMinComputeCapability = 60; +}; +struct Sm61 { + static int const kMinComputeCapability = 61; +}; + + +/// BIGISLAND Arch +struct Cu10 { + static int const kMinComputeCapability = 10; +}; + +struct Sm62 { + static int const kMinComputeCapability = 62; +}; + +/// Triggers a breakpoint on the device +CUTLASS_DEVICE +void device_breakpoint() { +#if defined(__CUDA_ARCH__) + asm volatile (" brkpt;\n"); +#endif + +//////////////////////////////////////////////////////////////////////////////////////////////////// +/// Switches to control performance improvement only supported on Iluvatar platform + +/// Compiler of Iluvatar-CoreX implicitly convert boolean type that is stored at VRF to a 64-bit +/// width integer type on SRF +#define IMPLICIT_VRF_BOOLEAN_TO_SRF_INTEGER 1 + +/// Enable block load or store +#define BLOCK_LOAD_STORE 1 + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +} +//////////////////////////////////////////////////////////////////////////////////////////////////// + +} // namespace arch +} // namespace cutlass + +//////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/cat_files/basic_gemm.cu b/cat_files/basic_gemm.cu new file mode 100644 index 00000000..bda012ab --- /dev/null +++ b/cat_files/basic_gemm.cu @@ -0,0 +1,492 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2020, 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. + * + **************************************************************************************************/ + +/* + This example demonstrates how to call a CUTLASS GEMM kernel and provides a naive reference + matrix multiply kernel to verify its correctness. + + The CUTLASS Gemm template is instantiated in the function CutlassSgemmNN. This is kernel computes + the general matrix product (GEMM) using single-precision floating-point arithmetic and assumes + all matrices have column-major layout. + + The threadblock tile size is chosen as 128x128x8 which offers good performance for large matrices. + See the CUTLASS Parallel for All blog post for more exposition on the tunable parameters available + in CUTLASS. + + https://devblogs.nvidia.com/cutlass-linear-algebra-cuda/ + + Aside from defining and launching the SGEMM kernel, this example does not use any other components + or utilities within CUTLASS. Such utilities are demonstrated elsewhere in other examples and are + prevalent in the CUTLASS unit tests. + + This example has delibrately been kept similar to the basic_gemm example from cutass-1.3 to + highlight the minimum amount of differences needed to transition to cutlass-2.0. + + Cutlass-1.3 sgemm: https://github.com/NVIDIA/cutlass/blob/master/examples/00_basic_gemm/basic_gemm.cu +*/ + +// Standard Library includes +#include +#include +#include + +// Helper methods to check for errors +#include "helper.h" + +// +// CUTLASS includes needed for single-precision GEMM kernel +// + +// Defines cutlass::gemm::device::Gemm, the generic Gemm computation template class. +#include "cutlass/gemm/device/gemm.h" + +/////////////////////////////////////////////////////////////////////////////////////////////////// +// +// This function defines a CUTLASS GEMM kernel instantiation, constructs its parameters object, +// and launches it on the CUDA device. +// +/////////////////////////////////////////////////////////////////////////////////////////////////// + +/// Define a CUTLASS GEMM template and launch a GEMM kernel. +cudaError_t CutlassSgemmNN( + int M, + int N, + int K, + float alpha, + float const *A, + int lda, + float const *B, + int ldb, + float beta, + float *C, + int ldc) { + + // Define type definition for single-precision CUTLASS GEMM with column-major + // input matrices and 128x128x8 threadblock tile size (chosen by default). + // + // To keep the interface manageable, several helpers are defined for plausible compositions + // including the following example for single-precision GEMM. Typical values are used as + // default template arguments. See `cutlass/gemm/device/default_gemm_configuration.h` for more details. + // + // To view the full gemm device API interface, see `cutlass/gemm/device/gemm.h` + + using ColumnMajor = cutlass::layout::ColumnMajor; + + using CutlassGemm = cutlass::gemm::device::Gemm; // Layout of C matrix + + // Define a CUTLASS GEMM type + CutlassGemm gemm_operator; + + // Construct the CUTLASS GEMM arguments object. + // + // One of CUTLASS's design patterns is to define gemm argument objects that are constructible + // in host code and passed to kernels by value. These may include pointers, strides, scalars, + // and other arguments needed by Gemm and its components. + // + // The benefits of this pattern are (1.) a structured, composable strategy for passing host-constructible + // arguments to kernels and (2.) minimized initialization overhead on kernel entry. + // + CutlassGemm::Arguments args({M , N, K}, // Gemm Problem dimensions + {A, lda}, // Tensor-ref for source matrix A + {B, ldb}, // Tensor-ref for source matrix B + {C, ldc}, // Tensor-ref for source matrix C + {C, ldc}, // Tensor-ref for destination matrix D (may be different memory than source C matrix) + {alpha, beta}); // Scalars used in the Epilogue + + // + // Launch the CUTLASS GEMM kernel. + // + + cutlass::Status status = gemm_operator(args); + + // + // Return a cudaError_t if the CUTLASS GEMM operator returned an error code. + // + + if (status != cutlass::Status::kSuccess) { + return cudaErrorUnknown; + } + + // Return success, if no errors were encountered. + return cudaSuccess; +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// +// +// The source code after this point in the file is generic CUDA using the CUDA Runtime API +// and simple CUDA kernels to initialize matrices and compute the general matrix product. +// +/////////////////////////////////////////////////////////////////////////////////////////////////// + +/// Kernel to initialize a matrix with small integers. +__global__ void InitializeMatrix_kernel( + float *matrix, + int ldm, + int rows, + int columns, + int seed = 0) { + + int i = threadIdx.x + blockIdx.x * blockDim.x; + int j = threadIdx.y + blockIdx.y * blockDim.y; + + if (i < rows && j < columns) { + int offset = i + j * ldm; + + // Generate arbitrary elements. + int const k = 16807; + int const m = 16; + float value = float(((offset + seed) * k % m) - m / 2); + + matrix[offset] = value; + } +} + +/// Simple function to initialize a matrix to arbitrary small integers. +cudaError_t InitializeMatrix(float *matrix, int ldm, int rows, int columns, int seed = 0) { + + dim3 block(16, 16); + dim3 grid( + (rows + block.x - 1) / block.x, + (columns + block.y - 1) / block.y + ); + + InitializeMatrix_kernel<<< grid, block >>>(matrix, ldm, rows, columns, seed); + + return cudaGetLastError(); +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// + +/// Allocates device memory for a matrix then fills with arbitrary small integers. +cudaError_t AllocateMatrix(float **matrix, int ldm, int rows, int columns, int seed = 0) { + cudaError_t result; + + size_t sizeof_matrix = sizeof(float) * ldm * columns; + + // Allocate device memory. + result = cudaMalloc(reinterpret_cast(matrix), sizeof_matrix); + + if (result != cudaSuccess) { + std::cerr << "Failed to allocate matrix: " + << cudaGetErrorString(result) << std::endl; + return result; + } + + // Clear the allocation. + result = cudaMemset(*matrix, 0, sizeof_matrix); + + if (result != cudaSuccess) { + std::cerr << "Failed to clear matrix device memory: " + << cudaGetErrorString(result) << std::endl; + return result; + } + + // Initialize matrix elements to arbitrary small integers. + result = InitializeMatrix(*matrix, ldm, rows, columns, seed); + + if (result != cudaSuccess) { + std::cerr << "Failed to initialize matrix: " + << cudaGetErrorString(result) << std::endl; + return result; + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// + +/// Naive reference GEMM computation. +__global__ void ReferenceGemm_kernel( + int M, + int N, + int K, + float alpha, + float const *A, + int lda, + float const *B, + int ldb, + float beta, + float *C, + int ldc) { + + int i = threadIdx.x + blockIdx.x * blockDim.x; + int j = threadIdx.y + blockIdx.y * blockDim.y; + + if (i < M && j < N) { + float accumulator = 0; + + for (int k = 0; k < K; ++k) { + accumulator += A[i + k * lda] * B[k + j * ldb]; + } + + C[i + j * ldc] = alpha * accumulator + beta * C[i + j * ldc]; + } +} + +/// Reference GEMM computation. +cudaError_t ReferenceGemm( + int M, + int N, + int K, + float alpha, + float const *A, + int lda, + float const *B, + int ldb, + float beta, + float *C, + int ldc) { + + dim3 block(16, 16); + dim3 grid( + (M + block.x - 1) / block.x, + (N + block.y - 1) / block.y + ); + + ReferenceGemm_kernel<<< grid, block >>>(M, N, K, alpha, A, lda, B, ldb, beta, C, ldc); + + return cudaGetLastError(); +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// + +/// Allocate several matrices in GPU device memory and call a single-precision +/// CUTLASS GEMM kernel. +cudaError_t TestCutlassGemm(int M, int N, int K, float alpha, float beta) { + cudaError_t result; + + // + // Define several matrices to be used as operands to GEMM kernels. + // + + // Compute leading dimensions for each matrix. + int lda = M; + int ldb = K; + int ldc = M; + + // Compute size in bytes of the C matrix. + size_t sizeof_C = sizeof(float) * ldc * N; + + // Define pointers to matrices in GPU device memory. + float *A; + float *B; + float *C_cutlass; + float *C_reference; + + // + // Allocate matrices in GPU device memory with arbitrary seeds. + // + + result = AllocateMatrix(&A, lda, M, K, 0); + + if (result != cudaSuccess) { + return result; + } + + result = AllocateMatrix(&B, ldb, K, N, 17); + + if (result != cudaSuccess) { + cudaFree(A); + return result; + } + + result = AllocateMatrix(&C_cutlass, ldc, M, N, 101); + + if (result != cudaSuccess) { + cudaFree(A); + cudaFree(B); + return result; + } + + result = AllocateMatrix(&C_reference, ldc, M, N, 101); + + if (result != cudaSuccess) { + cudaFree(A); + cudaFree(B); + cudaFree(C_cutlass); + return result; + } + + result = cudaMemcpy(C_reference, C_cutlass, sizeof_C, cudaMemcpyDeviceToDevice); + + if (result != cudaSuccess) { + std::cerr << "Failed to copy C_cutlass matrix to C_reference: " + << cudaGetErrorString(result) << std::endl; + + cudaFree(C_reference); + cudaFree(C_cutlass); + cudaFree(B); + cudaFree(A); + + return result; + } + + // + // Launch CUTLASS GEMM. + // + + result = CutlassSgemmNN(M, N, K, alpha, A, lda, B, ldb, beta, C_cutlass, ldc); + + if (result != cudaSuccess) { + std::cerr << "CUTLASS GEMM kernel failed: " + << cudaGetErrorString(result) << std::endl; + + cudaFree(C_reference); + cudaFree(C_cutlass); + cudaFree(B); + cudaFree(A); + + return result; + } + + // + // Verify. + // + + // Launch reference GEMM + result = ReferenceGemm(M, N, K, alpha, A, lda, B, ldb, beta, C_reference, ldc); + + if (result != cudaSuccess) { + std::cerr << "Reference GEMM kernel failed: " + << cudaGetErrorString(result) << std::endl; + + cudaFree(C_reference); + cudaFree(C_cutlass); + cudaFree(B); + cudaFree(A); + + return result; + } + + // Copy to host and verify equivalence. + std::vector host_cutlass(ldc * N, 0); + std::vector host_reference(ldc * N, 0); + + result = cudaMemcpy(host_cutlass.data(), C_cutlass, sizeof_C, cudaMemcpyDeviceToHost); + + if (result != cudaSuccess) { + std::cerr << "Failed to copy CUTLASS GEMM results: " + << cudaGetErrorString(result) << std::endl; + + cudaFree(C_reference); + cudaFree(C_cutlass); + cudaFree(B); + cudaFree(A); + + return result; + } + + result = cudaMemcpy(host_reference.data(), C_reference, sizeof_C, cudaMemcpyDeviceToHost); + + if (result != cudaSuccess) { + std::cerr << "Failed to copy Reference GEMM results: " + << cudaGetErrorString(result) << std::endl; + + cudaFree(C_reference); + cudaFree(C_cutlass); + cudaFree(B); + cudaFree(A); + + return result; + } + + // + // Free device memory allocations. + // + + cudaFree(C_reference); + cudaFree(C_cutlass); + cudaFree(B); + cudaFree(A); + + // + // Test for bit equivalence of results. + // + + if (host_cutlass != host_reference) { + std::cerr << "CUTLASS results incorrect." << std::endl; + + return cudaErrorUnknown; + } + + return cudaSuccess; +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// + +/// Entry point to basic_gemm example. +// +// usage: +// +// 00_basic_gemm +// +int main(int argc, const char *arg[]) { + + // + // Parse the command line to obtain GEMM dimensions and scalar values. + // + + // GEMM problem dimensions. + int problem[3] = { 128, 128, 128 }; + + for (int i = 1; i < argc && i < 4; ++i) { + std::stringstream ss(arg[i]); + ss >> problem[i - 1]; + } + + // Scalars used for linear scaling the result of the matrix product. + float scalars[2] = { 1, 0 }; + + for (int i = 4; i < argc && i < 6; ++i) { + std::stringstream ss(arg[i]); + ss >> scalars[i - 4]; + } + + // + // Run the CUTLASS GEMM test. + // + + cudaError_t result = TestCutlassGemm( + problem[0], // GEMM M dimension + problem[1], // GEMM N dimension + problem[2], // GEMM K dimension + scalars[0], // alpha + scalars[1] // beta + ); + + if (result == cudaSuccess) { + std::cout << "Passed." << std::endl; + } + + // Exit. + return result == cudaSuccess ? 0 : -1; +} + +/////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/cat_files/cutlass.h b/cat_files/cutlass.h new file mode 100644 index 00000000..3fb6f0a7 --- /dev/null +++ b/cat_files/cutlass.h @@ -0,0 +1,175 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2021, 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. + * + **************************************************************************************************/ + +/*************************************************************************************************** +* Copyright (c) 2021 Iluvatar CoreX. All rights reserved. +* Copyright Declaration: This software, including all of its code and documentation, +* except for the third-party software it contains, is a copyrighted work of Shanghai Iluvatar CoreX +* Semiconductor Co., Ltd. and its affiliates ("Iluvatar CoreX") in accordance with the PRC Copyright +* Law and relevant international treaties, and all rights contained therein are enjoyed by Iluvatar +* CoreX. No user of this software shall have any right, ownership or interest in this software and +* any use of this software shall be in compliance with the terms and conditions of the End User +* License Agreement. + **************************************************************************************************/ + +/*! \file + \brief Basic include for CUTLASS. +*/ + +#pragma once + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +#define CUTLASS_UNUSED(expr) do { (void)(expr); } while (0) + +#if defined(_MSC_VER) + #define CUTLASS_NOT_IMPLEMENTED() assert(0 && __FUNCSIG__) +#else + #define CUTLASS_NOT_IMPLEMENTED() assert(0 && __PRETTY_FUNCTION__) +#endif + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +namespace cutlass { + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__)) +#define CUTLASS_HOST_DEVICE __forceinline__ __device__ __host__ +#define CUTLASS_DEVICE __forceinline__ __device__ +#elif defined(__CUDACC_RTC__) +#define CUTLASS_HOST_DEVICE __forceinline__ __device__ +#define CUTLASS_DEVICE __forceinline__ __device__ +#else +#define CUTLASS_HOST_DEVICE inline +#define CUTLASS_DEVICE inline +#endif + +/// Status code returned by CUTLASS operations +enum class Status { + kSuccess, ///< Operation was successful. + kErrorMisalignedOperand, ///< operands fail alignment requirements. + kErrorInvalidDataType, ///< DataType fails requirement. + kErrorInvalidLayout, ///< Layout fails alignment requirement. + kErrorInvalidProblem, ///< Specified problem size is not supported by operator. + kErrorNotSupported, ///< Operation is not supported on current device. + kErrorWorkspaceNull, ///< The given workspace is null when it is required to be non-null. + kErrorInternal, ///< An error within CUTLASS occurred. + kErrorArchMismatch, ///< CUTLASS runs on a device that it was not compiled for. + kErrorInsufficientDriver, ///< CUTLASS runs with a driver that is too old. + kInvalid ///< Status is unspecified. +}; + +/// Convert cutlass status to status strings +CUTLASS_HOST_DEVICE +static char const* cutlassGetStatusString(cutlass::Status status) { + switch (status) { + case cutlass::Status::kSuccess: + return "Success"; + case cutlass::Status::kErrorMisalignedOperand: + return "Error Misaligned Operand"; + case cutlass::Status::kErrorInvalidDataType: + return "Error Invalid Data Type"; + case cutlass::Status::kErrorInvalidLayout: + return "Error Invalid Layout"; + case cutlass::Status::kErrorInvalidProblem: + return "Error Invalid Problem"; + case cutlass::Status::kErrorNotSupported: + return "Error Not Supported"; + case cutlass::Status::kErrorWorkspaceNull: + return "Error Workspace Null"; + case cutlass::Status::kErrorInternal: + return "Error Internal"; + case cutlass::Status::kErrorInsufficientDriver: + return "Error Insufficient Driver"; + case cutlass::Status::kErrorArchMismatch: + return "Erroor Architecture Mismatch"; + case cutlass::Status::kInvalid: break; + } + + return "Invalid status"; +} + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +#define CUTLASS_ASSERT(x) assert(x) + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +// CUTLASS_PRAGMA_(UNROLL|NO_UNROLL) optimization directives for the CUDA compiler. +#if defined(__CUDA_ARCH__) + #if defined(__CUDACC_RTC__) || (defined(__clang__) && defined(__CUDA__)) + #define CUTLASS_PRAGMA_UNROLL _Pragma("unroll") + #define CUTLASS_PRAGMA_NO_UNROLL _Pragma("unroll 1") + #else + #define CUTLASS_PRAGMA_UNROLL #pragma unroll + #define CUTLASS_PRAGMA_NO_UNROLL #pragma unroll 1 + #endif + + #define CUTLASS_GEMM_LOOP CUTLASS_PRAGMA_NO_UNROLL + +#else + + #define CUTLASS_PRAGMA_UNROLL + #define CUTLASS_PRAGMA_NO_UNROLL + #define CUTLASS_GEMM_LOOP + +#endif + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +static const int MEMORY_ACCESS_SIZE = 32; + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +static const int NUM_THREADS_PER_WARP = 64; +static const int NUM_THREADS_PER_HALF_WARP = NUM_THREADS_PER_WARP / 2; +static const int NUM_THREADS_PER_QUAD = 4; +static const int NUM_THREADS_PER_QUAD_PAIR = NUM_THREADS_PER_QUAD * 2; + +#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__)) + +/// Computes laneId within a warp +CUTLASS_DEVICE +int LaneId() { + return __ivcorex_lane_id(); +} + +/// Computes SM number the thread is running on +CUTLASS_DEVICE +int SmId() { + /// TODO(Peter Han): BI compiler doesn't support sm ID + __asm__ __volatile__("int3"); + return 0; +} + +#endif + +//////////////////////////////////////////////////////////////////////////////////////////////////// + +} // namespace cutlass + +//////////////////////////////////////////////////////////////////////////////////////////////////// + diff --git a/cat_files/gemm_batched_full.h b/cat_files/gemm_batched_full.h new file mode 100644 index 00000000..c37210e8 --- /dev/null +++ b/cat_files/gemm_batched_full.h @@ -0,0 +1,726 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2021, 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. + * + **************************************************************************************************/ + +/*************************************************************************************************** +* Copyright (c) 2021 Iluvatar CoreX. All rights reserved. +* Copyright Declaration: This software, including all of its code and documentation, +* except for the third-party software it contains, is a copyrighted work of Shanghai Iluvatar CoreX +* Semiconductor Co., Ltd. and its affiliates ("Iluvatar CoreX") in accordance with the PRC Copyright +* Law and relevant international treaties, and all rights contained therein are enjoyed by Iluvatar +* CoreX. No user of this software shall have any right, ownership or interest in this software and +* any use of this software shall be in compliance with the terms and conditions of the End User +* License Agreement. + **************************************************************************************************/ + +/*! \file + \brief Template for a pipelined GEMM kernel. Does not compute batching or support split-K. +*/ + +#pragma once + +#include "cutlass/cutlass.h" +#include "cutlass/numeric_types.h" +#include "cutlass/arch/arch.h" +#include "cutlass/device_kernel.h" + +#include "cutlass/gemm/threadblock/threadblock_swizzle.h" +#include "cutlass/gemm/kernel/gemm_batched.h" + +#include "cutlass/gemm/kernel/default_gemm.h" +#include "cutlass/gemm/device/default_gemm_configuration.h" + +//////////////////////////////////////////////////////////////////////////////// + +namespace cutlass { +namespace gemm { +namespace device { + +//////////////////////////////////////////////////////////////////////////////// + +/*! Gemm device-level operator. This is an interface to efficient CUTLASS GEMM kernels that may + be invoked from host code. + + The contributions of this class are: + + 1. At compile time, it maps data types and high-level structural parameters onto + specific CUTLASS components. + + 2. At runtime, it maps logical arguments to GEMM problems to kernel parameters. + + 3. At runtime, it launches kernels on the device. + + The intent is to provide a convenient mechanism for interacting with most plausible GEMM + configurations for each supported architecture. Consequently, not all parameters are exposed + to the top-level interface. Rather, sensible defaults at each level of the CUTLASS hierarchy + are selected to tradeoff simplicity of the interface with flexibility. We expect + most configurations to be specified at this level. Applications with more exotic requirements + may construct their kernels of interest using CUTLASS components at the threadblock, warp, + and thread levels of abstraction. + + CUTLASS exposes computations using the functor design pattern in which objects compose some + internal state with an overloaded function call operator. This enables decoupling of + initialization from execution, possibly reducing overhead during steady state phases of + application execution. + + CUTLASS device-level operators expose an Arguments structure encompassing each logical + input to the computation. This is distinct from the kernel-level Params structure pattern + which contains application-specific precomputed state needed by the device code. + + Example of a CUTLASS GEMM operator implementing the functionality of cuBLAS's SGEMM NN + is as follows: + + // + // Instantiate the CUTLASS GEMM operator. + // + + cutlass::gemm::device::Gemm< + float, + cutlass::layout::ColumnMajor, + float, + cutlass::layout::ColumnMajor, + float, + cutlass::layout::ColumnMajor + > gemm_op; + + // + // Launch the GEMM operation on the device + // + + cutlass::Status status = gemm_op({ + {m, n, k}, // GemmCoord problem_size, + {A, lda}, // TensorRef ref_A, + {B, ldb}, // TensorRef ref_B, + {C, ldc}, // TensorRef ref_C, + {D, ldd}, // TensorRef ref_D, + {alpha, beta} // EpilogueOutputOp::Params epilogue_op_params + }); + + + A simplified view of the template is listed below. + + template < + /// Element type for A matrix operand + typename ElementA, + + /// Layout type for A matrix operand + typename LayoutA, + + /// Element type for B matrix operand + typename ElementB, + + /// Layout type for B matrix operand + typename LayoutB, + + /// Element type for C and D matrix operands + typename ElementC, + + /// Layout type for C and D matrix operands + typename LayoutC, + + /// Element type for internal accumulation + typename ElementAccumulator, + + /// Operator class tag + typename OperatorClass, + + /// Tag indicating architecture to tune for. This is the minimum SM that + /// supports the intended feature. The device kernel can be built + /// targeting any SM larger than this number. + typename ArchTag, + + /// Threadblock-level tile size (concept: GemmShape) + typename ThreadblockShape, + + /// Warp-level tile size (concept: GemmShape) + typename WarpShape, + + /// Warp-level tile size (concept: GemmShape) + typename InstructionShape, + + /// Epilogue output operator + typename EpilogueOutputOp, + + /// Threadblock-level swizzling operator + typename ThreadblockSwizzle, + + /// Number of stages used in the pipelined mainloop + int Stages + > + class Gemm; +*/ +template < + /// Element type for A matrix operand + typename ElementA_, + /// Layout type for A matrix operand + typename LayoutA_, + /// Element type for B matrix operand + typename ElementB_, + /// Layout type for B matrix operand + typename LayoutB_, + /// Element type for C and D matrix operands + typename ElementC_, + /// Layout type for C and D matrix operands + typename LayoutC_, + /// Element type for internal accumulation + typename ElementAccumulator_ = ElementC_, + /// Operator class tag + typename OperatorClass_ = arch::OpClassSimt, + /// Tag indicating architecture to tune for + typename ArchTag_ = arch::Sm61, + /// Threadblock-level tile size (concept: GemmShape) + typename ThreadblockShape_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::ThreadblockShape, + /// Warp-level tile size (concept: GemmShape) + typename WarpShape_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::WarpShape, + /// Instruction-level tile size (concept: GemmShape) + typename InstructionShape_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::InstructionShape, + /// Epilogue output operator + typename EpilogueOutputOp_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::EpilogueOutputOp, + /// Threadblock-level swizzling operator + typename ThreadblockSwizzle_ = threadblock::GemmBatchedIdentityThreadblockSwizzle, + /// Number of stages used in the pipelined mainloop + int Stages = + DefaultGemmConfiguration::kStages, + /// Access granularity of A matrix in units of elements + int AlignmentA = + DefaultGemmConfiguration::kAlignmentA, + /// Access granularity of B matrix in units of elements + int AlignmentB = + DefaultGemmConfiguration::kAlignmentB, + /// Operation performed by GEMM + typename Operator_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::Operator +> +class GemmBatched { + public: + + using ElementA = ElementA_; + using LayoutA = LayoutA_; + using TensorRefA = TensorRef; + using ElementB = ElementB_; + using LayoutB = LayoutB_; + using TensorRefB = TensorRef; + using ElementC = ElementC_; + using LayoutC = LayoutC_; + using TensorRefC = TensorRef; + using TensorRefD = TensorRef; + using ElementAccumulator = ElementAccumulator_; + using OperatorClass = OperatorClass_; + using ArchTag = ArchTag_; + using ThreadblockShape = ThreadblockShape_; + using WarpShape = WarpShape_; + using InstructionShape = InstructionShape_; + using EpilogueOutputOp = EpilogueOutputOp_; + using ThreadblockSwizzle = ThreadblockSwizzle_; + static int const kStages = Stages; + static int const kAlignmentA = AlignmentA; + static int const kAlignmentB = AlignmentB; + static int const kAlignmentC = EpilogueOutputOp::kCount; + using Operator = Operator_; + + /// Define the kernel + using DefaultGemmKernel = typename kernel::DefaultGemm< + ElementA, + LayoutA, + kAlignmentA, + ElementB, + LayoutB, + kAlignmentB, + ElementC, + LayoutC, + ElementAccumulator, + OperatorClass, + ArchTag, + ThreadblockShape, + WarpShape, + InstructionShape, + EpilogueOutputOp, + ThreadblockSwizzle, + kStages, + false, + Operator + >::GemmKernel; + + using GemmKernel = kernel::GemmBatched; + + /// Argument structure + struct Arguments { + + // + // Data members + // + + GemmCoord problem_size; + TensorRef ref_A; + int64_t stride_A; + TensorRef ref_B; + int64_t stride_B; + TensorRef ref_C; + int64_t stride_C; + TensorRef ref_D; + int64_t stride_D; + typename EpilogueOutputOp::Params epilogue; + int batch_count; + + // + // Methods + // + + /// Default ctor + CUTLASS_HOST_DEVICE + Arguments() { } + + /// Constructs an Arguments structure + CUTLASS_HOST_DEVICE + Arguments( + GemmCoord problem_size_, + TensorRef ref_A_, + int64_t stride_A_, + TensorRef ref_B_, + int64_t stride_B_, + TensorRef ref_C_, + int64_t stride_C_, + TensorRef ref_D_, + int64_t stride_D_, + typename EpilogueOutputOp::Params epilogue_, + int batch_count_ + ): + problem_size(problem_size_), + ref_A(ref_A_), + stride_A(stride_A_), + ref_B(ref_B_), + stride_B(stride_B_), + ref_C(ref_C_), + stride_C(stride_C_), + ref_D(ref_D_), + stride_D(stride_D_), + epilogue(epilogue_), + batch_count(batch_count_) { } + }; + +private: + + /// Kernel parameters object + typename GemmKernel::Params params_; + +public: + + /// Constructs the GEMM. + GemmBatched() { } + + /// Determines whether the GEMM can execute the given problem. + static Status can_implement(Arguments const &args) { + + if (!TensorRef_aligned(args.ref_A, kAlignmentA) || (args.stride_A % kAlignmentA)) { + return Status::kErrorMisalignedOperand; + } + + if (!TensorRef_aligned(args.ref_B, kAlignmentB) || (args.stride_B % kAlignmentB)) { + return Status::kErrorMisalignedOperand; + } + + if (!TensorRef_aligned(args.ref_C, kAlignmentC) || (args.stride_C % kAlignmentC)) { + return Status::kErrorMisalignedOperand; + } + + if (!TensorRef_aligned(args.ref_D, kAlignmentC) || (args.stride_D % kAlignmentC)) { + return Status::kErrorMisalignedOperand; + } + + if ((args.problem_size.m() % kAlignmentA) || (args.problem_size.k() % kAlignmentA) || + (args.problem_size.n() % kAlignmentB) || (args.problem_size.k() % kAlignmentB) || + (args.problem_size.m() % kAlignmentC) || (args.problem_size.n() % kAlignmentC)) { + + return Status::kErrorMisalignedOperand; + } + + return Status::kSuccess; + } + + /// Gets the workspace size + static size_t get_workspace_size(Arguments const &args) { + return 0; + } + + /// Initializes GEMM state from arguments. + Status initialize(Arguments const &args, void *workspace = nullptr, cudaStream_t stream = nullptr) { + + // Determine grid shape + ThreadblockSwizzle threadblock_swizzle; + + cutlass::gemm::GemmCoord grid_shape = threadblock_swizzle.get_tiled_shape( + args.problem_size, + {ThreadblockShape::kM, ThreadblockShape::kN, ThreadblockShape::kK}, + args.batch_count); + + // Initialize the Params structure + params_ = typename GemmKernel::Params{ + args.problem_size, + grid_shape, + args.ref_A.non_const_ref(), + args.stride_A, + args.ref_B.non_const_ref(), + args.stride_B, + args.ref_C.non_const_ref(), + args.stride_C, + args.ref_D, + args.stride_D, + args.epilogue, + args.batch_count + }; + + return Status::kSuccess; + } + + /// Lightweight update given a subset of arguments + Status update(Arguments const &args, void *workspace = nullptr) { + + params_.ref_A.reset(args.ref_A.non_const_ref().data()); + params_.ref_B.reset(args.ref_B.non_const_ref().data()); + params_.ref_C.reset(args.ref_C.non_const_ref().data()); + params_.ref_D.reset(args.ref_D.data()); + + return Status::kSuccess; + } + + /// Runs the kernel using initialized state. + Status run(cudaStream_t stream = nullptr) { + + ThreadblockSwizzle threadblock_swizzle; + + dim3 grid = threadblock_swizzle.get_grid_shape(params_.grid_tiled_shape); + // XXX(Peter Han): prealod needs double warps in z direction + dim3 block(GemmKernel::kThreadCount, 1, kStages ? 1 : 2); + + cudaError_t result; + + int smem_size = int(sizeof(typename GemmKernel::SharedStorage)); + /// cudaFuncSetAttribute isn't supported under CUDA-8.0 + // if (smem_size >= (48 << 10)) { + // result = cudaFuncSetAttribute(Kernel, + // cudaFuncAttributeMaxDynamicSharedMemorySize, + // smem_size); + + // if (result != cudaSuccess) { + // return Status::kErrorInternal; + // } + + // result = cudaFuncSetAttribute( + // Kernel, + // cudaFuncAttributePreferredSharedMemoryCarveout, 100); + + // if (result != cudaSuccess) { + // return Status::kErrorInternal; + // } + // } + + cutlass::Kernel<<>>(params_); + + result = cudaGetLastError(); + + return result == cudaSuccess ? Status::kSuccess : Status::kErrorInternal; + } + + /// Runs the kernel using initialized state. + Status operator()(cudaStream_t stream = nullptr) { + return run(stream); + } + + /// Runs the kernel using initialized state. + Status operator()( + Arguments const &args, + void *workspace = nullptr, + cudaStream_t stream = nullptr) { + + Status status = initialize(args, workspace, stream); + + if (status == Status::kSuccess) { + status = run(stream); + } + + return status; + } +}; + +//////////////////////////////////////////////////////////////////////////////// + +/// Parital specialization for column-major output exchanges problem size and operand. +template < + /// Element type for A matrix operand + typename ElementA_, + /// Layout type for A matrix operand + typename LayoutA_, + /// Element type for B matrix operand + typename ElementB_, + /// Layout type for B matrix operand + typename LayoutB_, + /// Element type for C and D matrix operands + typename ElementC_, + /// Element type for internal accumulation + typename ElementAccumulator_, + /// Operator class tag + typename OperatorClass_, + /// Tag indicating architecture to tune for + typename ArchTag_, + /// Threadblock-level tile size (concept: GemmShape) + typename ThreadblockShape_, + /// Warp-level tile size (concept: GemmShape) + typename WarpShape_, + /// Warp-level tile size (concept: GemmShape) + typename InstructionShape_, + /// Epilogue output operator + typename EpilogueOutputOp_, + /// Threadblock-level swizzling operator + typename ThreadblockSwizzle_, + /// Number of stages used in the pipelined mainloop + int Stages, + /// Access granularity of A matrix in units of elements + int AlignmentA, + /// Access granularity of B matrix in units of elements + int AlignmentB, + typename Operator_ +> +class GemmBatched< + ElementA_, + LayoutA_, + ElementB_, + LayoutB_, + ElementC_, + layout::ColumnMajor, + ElementAccumulator_, + OperatorClass_, + ArchTag_, + ThreadblockShape_, + WarpShape_, + InstructionShape_, + EpilogueOutputOp_, + ThreadblockSwizzle_, + Stages, + AlignmentA, + AlignmentB, + Operator_ +> { +public: + + using ElementA = ElementA_; + using LayoutA = LayoutA_; + using TensorRefA = TensorRef; + using ElementB = ElementB_; + using LayoutB = LayoutB_; + using TensorRefB = TensorRef; + using ElementC = ElementC_; + using LayoutC = layout::ColumnMajor; + using TensorRefC = TensorRef; + using TensorRefD = TensorRef; + using ElementAccumulator = ElementAccumulator_; + using OperatorClass = OperatorClass_; + using ArchTag = ArchTag_; + using ThreadblockShape = ThreadblockShape_; + using WarpShape = WarpShape_; + using InstructionShape = InstructionShape_; + using EpilogueOutputOp = EpilogueOutputOp_; + using ThreadblockSwizzle = ThreadblockSwizzle_; + static int const kStages = Stages; + + static int const kAlignmentA = AlignmentA; + static int const kAlignmentB = AlignmentB; + static int const kAlignmentC = EpilogueOutputOp::kCount; + static bool const kSplitKSerial = false; + + // + using UnderlyingOperator = GemmBatched< + ElementB, + typename layout::LayoutTranspose::type, + ElementA, + typename layout::LayoutTranspose::type, + ElementC, + layout::RowMajor, + ElementAccumulator, + OperatorClass, + ArchTag, + ThreadblockShape, + WarpShape, + InstructionShape, + EpilogueOutputOp, + ThreadblockSwizzle, + Stages, + kAlignmentB, + kAlignmentA + >; + + using UnderlyingArguments = typename UnderlyingOperator::Arguments; + using GemmKernel = typename UnderlyingOperator::GemmKernel; + + /// Argument structure + struct Arguments { + + // + // Data members + // + + GemmCoord problem_size; + TensorRef ref_A; + int64_t stride_A; + TensorRef ref_B; + int64_t stride_B; + TensorRef ref_C; + int64_t stride_C; + TensorRef ref_D; + int64_t stride_D; + typename EpilogueOutputOp::Params epilogue; + int batch_count; + + // + // Methods + // + + /// Default ctor + CUTLASS_HOST_DEVICE + Arguments() { } + + /// Constructs an Arguments structure + CUTLASS_HOST_DEVICE + Arguments( + GemmCoord problem_size_, + TensorRef ref_A_, + int64_t stride_A_, + TensorRef ref_B_, + int64_t stride_B_, + TensorRef ref_C_, + int64_t stride_C_, + TensorRef ref_D_, + int64_t stride_D_, + typename EpilogueOutputOp::Params epilogue_, + int batch_count_ + ): + problem_size(problem_size_), + ref_A(ref_A_), + stride_A(stride_A_), + ref_B(ref_B_), + stride_B(stride_B_), + ref_C(ref_C_), + stride_C(stride_C_), + ref_D(ref_D_), + stride_D(stride_D_), + epilogue(epilogue_), + batch_count(batch_count_) { } + }; + +private: + + UnderlyingOperator underlying_operator_; + +public: + + /// Constructs the GEMM. + GemmBatched() { } + + /// Helper to construct a transposed equivalent for the underying GEMM operator + static UnderlyingArguments to_underlying_arguments(Arguments const &args) { + return UnderlyingArguments( + {args.problem_size.n(), args.problem_size.m(), args.problem_size.k()}, + {args.ref_B.data(), args.ref_B.stride(0)}, + args.stride_B, + {args.ref_A.data(), args.ref_A.stride(0)}, + args.stride_A, + {args.ref_C.data(), args.ref_C.stride(0)}, + args.stride_C, + {args.ref_D.data(), args.ref_D.stride(0)}, + args.stride_D, + args.epilogue, + args.batch_count + ); + } + + /// Determines whether the GEMM can execute the given problem. + static Status can_implement(Arguments const &args) { + + return UnderlyingOperator::can_implement(to_underlying_arguments(args)); + } + + /// Gets the workspace size + static size_t get_workspace_size(Arguments const &args) { + + return UnderlyingOperator::get_workspace_size(to_underlying_arguments(args)); + } + + /// Initializes GEMM state from arguments. + Status initialize(Arguments const &args, void *workspace = nullptr, cudaStream_t stream = nullptr) { + + return underlying_operator_.initialize(to_underlying_arguments(args), workspace); + } + + /// Lightweight update given a subset of arguments + Status update(Arguments const &args, void *workspace = nullptr) { + + return underlying_operator_.update(to_underlying_arguments(args), workspace); + } + + /// Runs the kernel using initialized state. + Status run(cudaStream_t stream = nullptr) { + + return underlying_operator_.run(stream); + } + + /// Runs the kernel using initialized state. + Status operator()(cudaStream_t stream = nullptr) { + return run(stream); + } + + /// Runs the kernel using initialized state. + Status operator()( + Arguments const &args, + void *workspace = nullptr, + cudaStream_t stream = nullptr) { + + Status status = initialize(args, workspace, stream); + + if (status == Status::kSuccess) { + status = run(stream); + } + + return status; + } + +}; + +//////////////////////////////////////////////////////////////////////////////// + +} // namespace device +} // namespace gemm +} // namespace cutlass + +//////////////////////////////////////////////////////////////////////////////// diff --git a/cat_files/gemm_device.h b/cat_files/gemm_device.h new file mode 100644 index 00000000..57fe147a --- /dev/null +++ b/cat_files/gemm_device.h @@ -0,0 +1,732 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2021, 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. + * + **************************************************************************************************/ + +/*************************************************************************************************** +* Copyright (c) 2021 Iluvatar CoreX. All rights reserved. +* Copyright Declaration: This software, including all of its code and documentation, +* except for the third-party software it contains, is a copyrighted work of Shanghai Iluvatar CoreX +* Semiconductor Co., Ltd. and its affiliates ("Iluvatar CoreX") in accordance with the PRC Copyright +* Law and relevant international treaties, and all rights contained therein are enjoyed by Iluvatar +* CoreX. No user of this software shall have any right, ownership or interest in this software and +* any use of this software shall be in compliance with the terms and conditions of the End User +* License Agreement. + **************************************************************************************************/ + +/*! \file + \brief Template for a pipelined GEMM kernel. Does not compute batching or support split-K. +*/ + +#pragma once + +#include "cutlass/cutlass.h" +#include "cutlass/numeric_types.h" +#include "cutlass/arch/arch.h" +#include "cutlass/device_kernel.h" + +#include "cutlass/gemm/threadblock/threadblock_swizzle.h" +#include "cutlass/gemm/kernel/gemm.h" + +#include "cutlass/gemm/kernel/default_gemm.h" +#include "cutlass/gemm/device/default_gemm_configuration.h" + +//////////////////////////////////////////////////////////////////////////////// + +namespace cutlass { +namespace gemm { +namespace device { + +///////////////////////////////////////////////////////////////////////////////////////////////// + +/*! Gemm device-level operator. This is an interface to efficient CUTLASS GEMM kernels that may + be invoked from host code. + + The contributions of this class are: + + 1. At compile time, it maps data types and high-level structural parameters onto + specific CUTLASS components. + + 2. At runtime, it maps logical arguments to GEMM problems to kernel parameters. + + 3. At runtime, it launches kernels on the device. + + The intent is to provide a convenient mechanism for interacting with most plausible GEMM + configurations for each supported architecture. Consequently, not all parameters are exposed + to the top-level interface. Rather, sensible defaults at each level of the CUTLASS hierarchy + are selected to tradeoff simplicity of the interface with flexibility. We expect + most configurations to be specified at this level. Applications with more exotic requirements + may construct their kernels of interest using CUTLASS components at the threadblock, warp, + and thread levels of abstraction. + + CUTLASS exposes computations using the functor design pattern in which objects compose some + internal state with an overloaded function call operator. This enables decoupling of + initialization from execution, possibly reducing overhead during steady state phases of + application execution. + + CUTLASS device-level operators expose an Arguments structure encompassing each logical + input to the computation. This is distinct from the kernel-level Params structure pattern + which contains application-specific precomputed state needed by the device code. + + Example of a CUTLASS GEMM operator implementing the functionality of cuBLAS's SGEMM NN + is as follows: + + // + // Instantiate the CUTLASS GEMM operator. + // + + cutlass::gemm::device::Gemm< + float, + cutlass::layout::ColumnMajor, + float, + cutlass::layout::ColumnMajor, + float, + cutlass::layout::ColumnMajor + > gemm_op; + + // + // Launch the GEMM operation on the device + // + + cutlass::Status status = gemm_op({ + {m, n, k}, // GemmCoord problem_size, + {A, lda}, // TensorRef ref_A, + {B, ldb}, // TensorRef ref_B, + {C, ldc}, // TensorRef ref_C, + {D, ldd}, // TensorRef ref_D, + {alpha, beta} // EpilogueOutputOp::Params epilogue_op_params + }); + + + A simplified view of the template is listed below. + + template < + /// Element type for A matrix operand + typename ElementA, + + /// Layout type for A matrix operand + typename LayoutA, + + /// Element type for B matrix operand + typename ElementB, + + /// Layout type for B matrix operand + typename LayoutB, + + /// Element type for C and D matrix operands + typename ElementC, + + /// Layout type for C and D matrix operands + typename LayoutC, + + /// Element type for internal accumulation + typename ElementAccumulator, + + /// Operator class tag + typename OperatorClass, + + /// Tag indicating architecture to tune for. This is the minimum SM that + /// supports the intended feature. The device kernel can be built + /// targeting any SM larger than this number. + typename ArchTag, + + /// Threadblock-level tile size (concept: GemmShape) + typename ThreadblockShape, + + /// Warp-level tile size (concept: GemmShape) + typename WarpShape, + + /// Warp-level tile size (concept: GemmShape) + typename InstructionShape, + + /// Epilogue output operator + typename EpilogueOutputOp, + + /// Threadblock-level swizzling operator + typename ThreadblockSwizzle, + + /// Number of stages used in the pipelined mainloop + int Stages + > + class Gemm; +*/ +template < + /// Element type for A matrix operand + typename ElementA_, + /// Layout type for A matrix operand + typename LayoutA_, + /// Element type for B matrix operand + typename ElementB_, + /// Layout type for B matrix operand + typename LayoutB_, + /// Element type for C and D matrix operands + typename ElementC_, + /// Layout type for C and D matrix operands + typename LayoutC_, + /// Element type for internal accumulation + typename ElementAccumulator_ = ElementC_, + /// Operator class tag + typename OperatorClass_ = arch::OpClassSimt, + /// Tag indicating architecture to tune for + typename ArchTag_ = arch::Sm61, + /// Threadblock-level tile size (concept: GemmShape) + typename ThreadblockShape_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::ThreadblockShape, + /// Warp-level tile size (concept: GemmShape) + typename WarpShape_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::WarpShape, + /// Instruction-level tile size (concept: GemmShape) + typename InstructionShape_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::InstructionShape, + /// Epilogue output operator + typename EpilogueOutputOp_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::EpilogueOutputOp, + /// Threadblock-level swizzling operator + typename ThreadblockSwizzle_ = + typename threadblock::GemmIdentityThreadblockSwizzle<>, + /// Number of stages used in the pipelined mainloop + int Stages = + DefaultGemmConfiguration::kStages, + /// Access granularity of A matrix in units of elements + int AlignmentA = + DefaultGemmConfiguration::kAlignmentA, + /// Access granularity of B matrix in units of elements + int AlignmentB = + DefaultGemmConfiguration::kAlignmentB, + /// If true, kernel supports split-K with serial reduction + bool SplitKSerial = false, + /// Operation performed by GEMM + typename Operator_ = typename DefaultGemmConfiguration< + OperatorClass_, ArchTag_, ElementA_, ElementB_, ElementC_, + ElementAccumulator_>::Operator> +class Gemm { + public: + + using ElementA = ElementA_; + using LayoutA = LayoutA_; + using TensorRefA = TensorRef; + using ElementB = ElementB_; + using LayoutB = LayoutB_; + using TensorRefB = TensorRef; + using ElementC = ElementC_; + using LayoutC = LayoutC_; + using TensorRefC = TensorRef; + using TensorRefD = TensorRef; + using ElementAccumulator = ElementAccumulator_; + using OperatorClass = OperatorClass_; + using ArchTag = ArchTag_; + using ThreadblockShape = ThreadblockShape_; + using WarpShape = WarpShape_; + using InstructionShape = InstructionShape_; + using EpilogueOutputOp = EpilogueOutputOp_; + using ThreadblockSwizzle = ThreadblockSwizzle_; + using Operator = Operator_; + static int const kStages = Stages; + static int const kAlignmentA = AlignmentA; + static int const kAlignmentB = AlignmentB; + static int const kAlignmentC = EpilogueOutputOp::kCount; + static bool const kSplitKSerial = SplitKSerial; + static ComplexTransform const kTransformA = ComplexTransform::kNone; + static ComplexTransform const kTransformB = ComplexTransform::kNone; + + /// Define the kernel + using GemmKernel = typename kernel::DefaultGemm< + ElementA, + LayoutA, + kAlignmentA, + ElementB, + LayoutB, + kAlignmentB, + ElementC, + LayoutC, + ElementAccumulator, + OperatorClass, + ArchTag, + ThreadblockShape, + WarpShape, + InstructionShape, + EpilogueOutputOp, + ThreadblockSwizzle, + kStages, + kSplitKSerial, + Operator + >::GemmKernel; + + /// Argument structure + struct Arguments { + + // + // Data members + // + + GemmCoord problem_size; + TensorRef ref_A; + TensorRef ref_B; + TensorRef ref_C; + TensorRef ref_D; + typename EpilogueOutputOp::Params epilogue; + int split_k_slices; + + // + // Methods + // + + /// Default ctor + CUTLASS_HOST_DEVICE + Arguments(): problem_size(0, 0, 0), split_k_slices(1) { + + } + + /// Constructs an Arguments structure + CUTLASS_HOST_DEVICE + Arguments( + GemmCoord problem_size_, + TensorRef ref_A_, + TensorRef ref_B_, + TensorRef ref_C_, + TensorRef ref_D_, + typename EpilogueOutputOp::Params epilogue_ = + typename EpilogueOutputOp::Params(), + int split_k_slices = 1 + ): + problem_size(problem_size_), + ref_A(ref_A_), + ref_B(ref_B_), + ref_C(ref_C_), + ref_D(ref_D_), + epilogue(epilogue_), + split_k_slices(split_k_slices) { + + } + }; + +private: + + /// Kernel parameters object + typename GemmKernel::Params params_; + +public: + + /// Constructs the GEMM. + Gemm() { } + + /// Determines whether the GEMM can execute the given problem. + static Status can_implement(Arguments const &args) { + + if (!kSplitKSerial && args.split_k_slices > 1) { + return Status::kErrorInvalidProblem; + } + + Status status = GemmKernel::can_implement( + args.problem_size, + args.ref_A.non_const_ref(), + args.ref_B.non_const_ref(), + args.ref_C.non_const_ref(), + args.ref_D + ); + + if (status != Status::kSuccess) { + return status; + } + + return Status::kSuccess; + } + + /// Gets the workspace size + static size_t get_workspace_size(Arguments const &args) { + + size_t bytes = 0; + + // Determine grid shape + ThreadblockSwizzle threadblock_swizzle; + + cutlass::gemm::GemmCoord tiled_shape = threadblock_swizzle.get_tiled_shape( + args.problem_size, + {ThreadblockShape::kM, ThreadblockShape::kN, ThreadblockShape::kK}, + args.split_k_slices); + + if (kSplitKSerial && args.split_k_slices > 1) { + + bytes += sizeof(int) * size_t(tiled_shape.m()) * size_t(tiled_shape.n()); + } + + return bytes; + } + + /// Initializes GEMM state from arguments. + Status initialize(Arguments const &args, void *workspace = nullptr, cudaStream_t stream = nullptr) { + + // Determine grid shape + ThreadblockSwizzle threadblock_swizzle; + + cutlass::gemm::GemmCoord grid_shape = threadblock_swizzle.get_tiled_shape( + args.problem_size, + {ThreadblockShape::kM, ThreadblockShape::kN, ThreadblockShape::kK}, + args.split_k_slices); + + if (kSplitKSerial) { + if (args.split_k_slices > 1) { + if (!workspace) { + return Status::kErrorWorkspaceNull; + } + + size_t bytes = get_workspace_size(args); + + cudaError_t result = cudaMemsetAsync(workspace, 0, bytes, stream); + + if (result != cudaSuccess) { + return Status::kErrorInternal; + } + } + } + else { + + if (args.split_k_slices > 1) { + return Status::kErrorInvalidProblem; + } + } + + // Initialize the Params structure + params_ = typename GemmKernel::Params{ + args.problem_size, + grid_shape, + args.ref_A.non_const_ref(), + args.ref_B.non_const_ref(), + args.ref_C.non_const_ref(), + args.ref_D, + args.epilogue, + static_cast(workspace) + }; + + return Status::kSuccess; + } + + /// Lightweight update given a subset of arguments + Status update(Arguments const &args, void *workspace = nullptr) { + + if (kSplitKSerial && args.split_k_slices > 1) { + if (!workspace) { + return Status::kErrorWorkspaceNull; + } + } + + params_.ref_A.reset(args.ref_A.non_const_ref().data()); + params_.ref_B.reset(args.ref_B.non_const_ref().data()); + params_.ref_C.reset(args.ref_C.non_const_ref().data()); + params_.ref_D.reset(args.ref_D.data()); + params_.output_op = args.epilogue; + params_.semaphore = static_cast(workspace); + + return Status::kSuccess; + } + + /// Runs the kernel using initialized state. + Status run(cudaStream_t stream = nullptr) { + + ThreadblockSwizzle threadblock_swizzle; + + dim3 grid = threadblock_swizzle.get_grid_shape(params_.grid_tiled_shape); + // XXX(Peter Han): prealod needs double warps in z direction + dim3 block(GemmKernel::kThreadCount, 1, kStages ? 1 : 2); + + cudaError_t result; + + int smem_size = int(sizeof(typename GemmKernel::SharedStorage)); + /// cudaFuncSetAttribute isn't supported under CUDA-8.0 + // if (smem_size >= (48 << 10)) { + // result = cudaFuncSetAttribute(Kernel, + // cudaFuncAttributeMaxDynamicSharedMemorySize, + // smem_size); + + // if (result != cudaSuccess) { + // return Status::kErrorInternal; + // } + + // result = cudaFuncSetAttribute( + // Kernel, + // cudaFuncAttributePreferredSharedMemoryCarveout, 100); + + // if (result != cudaSuccess) { + // return Status::kErrorInternal; + // } + // } + + cutlass::Kernel<<>>(params_); + + result = cudaGetLastError(); + + return result == cudaSuccess ? Status::kSuccess : Status::kErrorInternal; + } + + /// Runs the kernel using initialized state. + Status operator()(cudaStream_t stream = nullptr) { + return run(stream); + } + + /// Runs the kernel using initialized state. + Status operator()( + Arguments const &args, + void *workspace = nullptr, + cudaStream_t stream = nullptr) { + + Status status = initialize(args, workspace, stream); + + if (status == Status::kSuccess) { + status = run(stream); + } + + return status; + } +}; + +//////////////////////////////////////////////////////////////////////////////// + +/// Parital specialization for column-major output exchanges problem size and operand. +template < + /// Element type for A matrix operand + typename ElementA_, + /// Layout type for A matrix operand + typename LayoutA_, + /// Element type for B matrix operand + typename ElementB_, + /// Layout type for B matrix operand + typename LayoutB_, + /// Element type for C and D matrix operands + typename ElementC_, + /// Element type for internal accumulation + typename ElementAccumulator_, + /// Operator class tag + typename OperatorClass_, + /// Tag indicating architecture to tune for + typename ArchTag_, + /// Threadblock-level tile size (concept: GemmShape) + typename ThreadblockShape_, + /// Warp-level tile size (concept: GemmShape) + typename WarpShape_, + /// Instruction-level tile size (concept: GemmShape) + typename InstructionShape_, + /// Epilogue output operator + typename EpilogueOutputOp_, + /// Threadblock-level swizzling operator + typename ThreadblockSwizzle_, + /// Number of stages used in the pipelined mainloop + int Stages, + /// Access granularity of A matrix in units of elements + int AlignmentA, + /// Access granularity of B matrix in units of elements + int AlignmentB, + /// If true, kernel supports split-K as a serial reduction + bool SplitKSerial, + /// Operation performed by GEMM + typename Operator_> +class Gemm { + public: + + using ElementA = ElementA_; + using LayoutA = LayoutA_; + using TensorRefA = TensorRef; + using ElementB = ElementB_; + using LayoutB = LayoutB_; + using TensorRefB = TensorRef; + using ElementC = ElementC_; + using LayoutC = layout::ColumnMajor; + using TensorRefC = TensorRef; + using TensorRefD = TensorRef; + using ElementAccumulator = ElementAccumulator_; + using OperatorClass = OperatorClass_; + using ArchTag = ArchTag_; + using ThreadblockShape = ThreadblockShape_; + using WarpShape = WarpShape_; + using InstructionShape = InstructionShape_; + using EpilogueOutputOp = EpilogueOutputOp_; + using ThreadblockSwizzle = ThreadblockSwizzle_; + using Operator = Operator_; + static int const kStages = Stages; + static int const kAlignmentA = AlignmentA; + static int const kAlignmentB = AlignmentB; + static ComplexTransform const kTransformA = ComplexTransform::kNone; + static ComplexTransform const kTransformB = ComplexTransform::kNone; + static bool const kSplitKSerial = SplitKSerial; + + using UnderlyingOperator = Gemm< + ElementB, + typename layout::LayoutTranspose::type, + ElementA, + typename layout::LayoutTranspose::type, + ElementC, + layout::RowMajor, + ElementAccumulator, + OperatorClass, + ArchTag, + ThreadblockShape, + WarpShape, + InstructionShape, + EpilogueOutputOp, + ThreadblockSwizzle, + Stages, + kAlignmentB, + kAlignmentA, + SplitKSerial, + Operator + >; + + using UnderlyingArguments = typename UnderlyingOperator::Arguments; + using GemmKernel = typename UnderlyingOperator::GemmKernel; + static int const kAlignmentC = UnderlyingOperator::kAlignmentC; + + /// Argument structure + struct Arguments { + + // + // Data members + // + + GemmCoord problem_size; + TensorRef ref_A; + TensorRef ref_B; + TensorRef ref_C; + TensorRef ref_D; + typename EpilogueOutputOp::Params epilogue; + int split_k_slices; + + // + // Methods + // + + /// Default ctor + CUTLASS_HOST_DEVICE + Arguments() { } + + /// Constructs an Arguments structure + CUTLASS_HOST_DEVICE + Arguments( + GemmCoord problem_size_, + TensorRef ref_A_, + TensorRef ref_B_, + TensorRef ref_C_, + TensorRef ref_D_, + typename EpilogueOutputOp::Params epilogue_ = + typename EpilogueOutputOp::Params(), + int split_k_slices = 1 + ): + problem_size(problem_size_), + ref_A(ref_A_), + ref_B(ref_B_), + ref_C(ref_C_), + ref_D(ref_D_), + epilogue(epilogue_), + split_k_slices(split_k_slices) { } + }; + +private: + + UnderlyingOperator underlying_operator_; + +public: + + /// Constructs the GEMM. + Gemm() { } + + /// Helper to construct a transposed equivalent for the underying GEMM operator + static UnderlyingArguments to_underlying_arguments(Arguments const &args) { + return UnderlyingArguments( + {args.problem_size.n(), args.problem_size.m(), args.problem_size.k()}, + {args.ref_B.data(), args.ref_B.stride(0)}, + {args.ref_A.data(), args.ref_A.stride(0)}, + {args.ref_C.data(), args.ref_C.stride(0)}, + {args.ref_D.data(), args.ref_D.stride(0)}, + args.epilogue, + args.split_k_slices + ); + } + + /// Determines whether the GEMM can execute the given problem. + static Status can_implement(Arguments const &args) { + + return UnderlyingOperator::can_implement(to_underlying_arguments(args)); + } + + /// Gets the workspace size + static size_t get_workspace_size(Arguments const &args) { + + return UnderlyingOperator::get_workspace_size(to_underlying_arguments(args)); + } + + /// Initializes GEMM state from arguments. + Status initialize(Arguments const &args, void *workspace = nullptr, cudaStream_t stream = nullptr) { + + return underlying_operator_.initialize(to_underlying_arguments(args), workspace, stream); + } + + /// Lightweight update given a subset of arguments + Status update(Arguments const &args, void *workspace = nullptr) { + + return underlying_operator_.update(to_underlying_arguments(args), workspace); + } + + /// Runs the kernel using initialized state. + Status run(cudaStream_t stream = nullptr) { + + return underlying_operator_.run(stream); + } + + /// Runs the kernel using initialized state. + Status operator()(cudaStream_t stream = nullptr) { + return run(stream); + } + + /// Runs the kernel using initialized state. + Status operator()( + Arguments const &args, + void *workspace = nullptr, + cudaStream_t stream = nullptr) { + + Status status = initialize(args, workspace, stream); + + if (status == Status::kSuccess) { + status = run(stream); + } + + return status; + } +}; + +//////////////////////////////////////////////////////////////////////////////// + +} // namespace device +} // namespace gemm +} // namespace cutlass + +//////////////////////////////////////////////////////////////////////////////// diff --git a/cat_files/turing_tensorop_gemm.cu b/cat_files/turing_tensorop_gemm.cu new file mode 100644 index 00000000..d18a4e6a --- /dev/null +++ b/cat_files/turing_tensorop_gemm.cu @@ -0,0 +1,354 @@ +/*************************************************************************************************** + * Copyright (c) 2017-2020, 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. + * + **************************************************************************************************/ + +/** +This example shows how to run matrix multiplication kernels using functions and data structures +provided by CUTLASS using tensor cores; which we run on a NVIDIA Turing GPU. + +Writing a single high performance matrix multiplication kernel is hard but do-able. Whereas writing +high performance kernels at scale which works for multiple problem sizes with good abstractions is +really hard. CUTLASS solves this problem by providing simplified abstractions to compose +multiple sections of gemm kernel. When used properly, the kernels can hit peak performance of GPU +easily. + +CUTLASS divides a kernel into hierarchical composable sections. Which means, at each thread, warp +and thread-block level, they compute on their own tile-size with higher level of tile sizes being +composed from lower level ones. Multiple thread-tiles (tile size each thread computes) can be used +to form warp-tiles (tile size each warp computes) and multiple warp tiles can be used to compute +threadblock-tile (tile size computed by a threadblock). + +In thie example, we split variable initialization into +1. Setting up data properties : describes how matrices are laid out in the memory and how the kernel +can view them (logical to physical mapping) +2. Setting up computation properties : describes how the above set matrices will be used to compute +output of matrix multiplication. + +First, we setup the data types of matrices A, B, C and D along with alpha, beta as the equation for +GEMM is D = alpha * A * B + beta * C. In CUTLASS, the kernels first compute A * B and leaves the +rest of the computation to end of the kernel as alpha * X + beta * C is a simple element-wise +operation on X (A * B) and C. We call this as epilogue of kernel. Hence, we setup data types for +alpha and beta to be equal to ElementComputeEpilogue = int32_t. As we want to use MMA instructions +on Turing and they support 8-bit signed integer (int8_t), we use data type for elements in input +matrix A and B as int8_t. Volta also supports accumulation of partial dot product to int32_t, which +can store wider range of numbers, we use it as data type of output matrix elements and accumulation. +We convey this to CUTLASS kernel by initializing template variables ElementAccumulator (int32_t), +ElementComputeEpilogue (int32_t), ElementInputA (int8_t), ElementInputB (int8_t), ElementOutput +(int32_t). Communicating just the data type is not enough. As the data is laid out linearly in +memory, we have to convey the layout of matrices. We do that by initializing template variable +LayoutInputA to column major cutlass variable, LayoutInputB to row major and LayoutOutput to row +major. Next, we setup rules to comptue alpha * X + beta * C which is called epilogue of the kernel. +We initialize template variable EpilogueOp, which takes the data type of output ElementOutput +(int32_t), the number of elements per vector memory access (16), data type of accumulator (int32_t) +and data type of computation of linear combination (alpha * X + beta * C). + +Now that we setup the properties of data, we have to setup properties of computation. + +Second, we create template variables of tile sizes for thread-block, warp and mma-op to 128x256x64, +64x64x16, 8x8x16 (MxNxK) respectively. When passed to instantiate CUTLASS GEMM kernel, it internally +deduce the amount of threads needed per thread-block, amount of shared memory, storing data in +bank-conflict free manner, and ton of other variables required to compose, intialize and launch a +high performance GEMM kernel. This is the beauty of CUTLASS, it relieves developer from +understanding and coding complicated hardware optimizations which can easily go wrong. + +CUTLASS also supports multiple MMA pipelines in a threadblock. What are MMA pipelines? MMA pipelines +constitute the whole process of loading input data from global memory to shared memory, loading data +from shared memory to registers, doing matrix multiplication, store to global memory. The below flow +sequence shows a typical mma pipeline. + +matrix in global memory -> registers -> tile in shared memory -> registers -> mma -> registers -> +output to global memory + +The problem with single pipeline is, each stage is synchronous which means, each stage has to wait +until the previous finished executing. There are stages in the pipeline which do not have fixed +latency, for example, the loads from global memory and shared memory. Therefore, we can add one more +pipeline with a phase shift in mma kernel to hide latency from global and shared memory loads. +Finally, the pipeline in a kernel looks like + +(1) matrix in global memory -> (2) registers -> (3) tile in shared memory -> (4) registers -> (5) +mma -> (6) registers -> (7) output to global memory (1) -> (2) -> (3) matrix in global +memory -> (4) registers -> (5) tile in shared memory -> (6) registers -> (7) mma -> (8) registers -> +(9) output to global memory + +This way, you can hide the second global memoroy load latency by doing computation on already loaded +input data. + +There are few more template variables initialized such as, which threadblock tile of output matrix +is done which threadblock launched on an SM, CUDA SM architecture of GPU you want to run on. + +These are all put together to create a template variable which describes CUTLASS GEMM kernel using +cutlass::gemm::device::Gemm template. + +The next step is to intialize physical data, instantiate and initialize CUTLASS kernel and run it. +We use CUTLASS utilities to initialize, fill, compare matrices as they are simple and doesn't come +in the way of learning CUTLASS. + +Once all the matrices are initialized and filled with data, create arguments tuple to launch CUTLASS +kernel which takes problem size (M = 5120, N = 4096 and K = 4096), matrices, alpha, beta and the +important one, split k-dimension factor. Along with that, we query CUTLASS if any scratch-space +memory required by the kernel we instantiated. If yes, we create it and pass it along with other +arguments created to intialize CUTLASS kernel then, the kernel is launched. + +In this example, we later on launch a reference gemm kernel (from CUTLASS utilities) to compare if +the output from CUTLASS kernel is same as reference GEMM kernel. +*/ + +#include + +#include "cutlass/cutlass.h" +#include "cutlass/gemm/device/gemm.h" +#include "cutlass/util/host_tensor.h" +#include "cutlass/util/reference/device/gemm.h" +#include "cutlass/util/reference/host/tensor_compare.h" +#include "cutlass/util/reference/host/tensor_copy.h" +#include "cutlass/util/reference/host/tensor_fill.h" +#include "cutlass/util/tensor_view_io.h" +#include "helper.h" + +// The code section below describes datatype for input, output matrices and computation between +// elements in input matrices. +using ElementAccumulator = int32_t; // <- data type of accumulator +using ElementComputeEpilogue = ElementAccumulator; // <- data type of epilogue operations +using ElementInputA = int8_t; // <- data type of elements in input matrix A +using ElementInputB = int8_t; // <- data type of elements in input matrix B +using ElementOutput = int32_t; // <- data type of elements in output matrix D + +// The code section below describes matrix layout of input and output matrices. Column Major for +// Matrix A, Row Major for Matrix B and Row Major for Matrix C +using LayoutInputA = cutlass::layout::RowMajor; +using LayoutInputB = cutlass::layout::ColumnMajor; +using LayoutOutput = cutlass::layout::RowMajor; + +// This code section describes whether you want to use tensor cores or regular SIMT cores on GPU SM +using MMAOp = cutlass::arch::OpClassTensorOp; + +// This code section describes CUDA SM architecture number +using SmArch = cutlass::arch::Sm75; + +// This code section describes the tile size a thread block will compute +using ShapeMMAThreadBlock = + cutlass::gemm::GemmShape<128, 256, 64>; // <- threadblock tile M = 128, N = 256, K = 64 +// This code section describes tile size a warp will compute +using ShapeMMAWarp = cutlass::gemm::GemmShape<64, 64, 64>; // <- warp tile M = 64, N = 64, K = 64 +// This code section describes the size of MMA op +using ShapeMMAOp = cutlass::gemm::GemmShape<8, 8, 16>; // <- MMA Op tile M = 8, N = 8, K = 16 + +// This code section describes how threadblocks are scheduled on GPU +using SwizzleThreadBlock = cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>; // <- ?? + +// This code section describes the epilogue part of the kernel +using EpilogueOp = cutlass::epilogue::thread::LinearCombination< + ElementOutput, // <- data type of output matrix + 128 / cutlass::sizeof_bits::value, // <- the number of elements per vectorized + // memory access. For a byte, it's 16 + // elements. This becomes the vector width of + // math instructions in the epilogue too + ElementAccumulator, // <- data type of accumulator + ElementComputeEpilogue>; // <- data type for alpha/beta in linear combination function + +// Number of pipelines you want to use +constexpr int NumStages = 2; + +using Gemm = cutlass::gemm::device::Gemm; + +int run() { + + // Turing Tensor Core operations exposed with mma.sync and ldmatrix are first available + // in CUDA 10.2. + // + // CUTLASS must be compiled with CUDA 10.2 Toolkit to run these examples. + if (!(__CUDACC_VER_MAJOR__ > 10 || (__CUDACC_VER_MAJOR__ == 10 && __CUDACC_VER_MINOR__ >= 2))) { + std::cerr << "Turing Tensor Core operations must be compiled with CUDA 10.2 Toolkit or later." << std::endl; + return -1; + } + + cudaDeviceProp props; + + cudaError_t error = cudaGetDeviceProperties(&props, 0); + if (error != cudaSuccess) { + std::cerr << "cudaGetDeviceProperties() returned an error: " << cudaGetErrorString(error) << std::endl; + return -1; + } + + if (!((props.major * 10 + props.minor) >= 75)) { + std::cerr << "Turing Tensor Core operations must be run on a machine with compute capability at least 75." + << std::endl; + + // Return 0 so tests are considered passing if run on unsupported platforms. + return 0; + } + + const int length_m = 5120; + const int length_n = 4096; + const int length_k = 4096; + + // Create a tuple of problem size for matrix multiplication + cutlass::gemm::GemmCoord problem_size(length_m, length_n, length_k); + + // Initialize tensors using CUTLASS helper functions + cutlass::HostTensor tensor_a( + problem_size.mk()); // <- Create matrix A with dimensions M x K + cutlass::HostTensor tensor_b( + problem_size.kn()); // <- Create matrix B with dimensions K x N + cutlass::HostTensor tensor_c( + problem_size.mn()); // <- Create matrix C with dimensions M x N + cutlass::HostTensor tensor_d( + problem_size.mn()); // <- Create matrix D with dimensions M x N used to store output from + // CUTLASS kernel + cutlass::HostTensor tensor_ref_d( + problem_size.mn()); // <- Create matrix D with dimensions M x N used to store output from + // reference kernel + + // Fill input and output matrices on host using CUTLASS helper functions + cutlass::reference::host::TensorFillRandomUniform( + tensor_a.host_view(), + 1, + ElementInputA(4), + ElementInputA(-4), + 0); // <- Fill matrix A on host with uniform-distribution random data + cutlass::reference::host::TensorFillRandomUniform( + tensor_b.host_view(), + 1, + ElementInputB(4), + ElementInputB(-4), + 0); // <- Fill matrix B on host with uniform-distribution random data + cutlass::reference::host::TensorFillRandomUniform( + tensor_c.host_view(), + 1, + ElementOutput(4), + ElementOutput(-4), + 0); // <- Fill matrix C on host with uniform-distribution random data + cutlass::reference::host::TensorFill( + tensor_d.host_view()); // <- fill matrix D on host with zeros + cutlass::reference::host::TensorFill( + tensor_ref_d.host_view()); // <- fill matrix D for reference on host with zeros + + // Copy data from host to GPU + tensor_a.sync_device(); + tensor_b.sync_device(); + tensor_c.sync_device(); + tensor_d.sync_device(); + tensor_ref_d.sync_device(); + + // Initialize alpha and beta for dot product computation + ElementComputeEpilogue alpha = ElementComputeEpilogue(1); + ElementComputeEpilogue beta = ElementComputeEpilogue(0); + + // Split K dimension into 1 partitions + int split_k_slices = 1; + + // Create a tuple of gemm kernel arguments. This is later passed as arguments to launch + // instantiated CUTLASS kernel + typename Gemm::Arguments arguments{problem_size, // <- problem size of matrix multiplication + tensor_a.device_ref(), // <- reference to matrix A on device + tensor_b.device_ref(), // <- reference to matrix B on device + tensor_c.device_ref(), // <- reference to matrix C on device + tensor_d.device_ref(), // <- reference to matrix D on device + {alpha, beta}, // <- tuple of alpha and beta + split_k_slices}; // <- k-dimension split factor + + // Using the arguments, query for extra workspace required for matrix multiplication computation + size_t workspace_size = Gemm::get_workspace_size(arguments); + + // Allocate workspace memory + cutlass::device_memory::allocation workspace(workspace_size); + + // Instantiate CUTLASS kernel depending on templates + Gemm gemm_op; + + // Initialize CUTLASS kernel with arguments and workspace pointer + cutlass::Status status = gemm_op.initialize(arguments, workspace.get()); + CUTLASS_CHECK(status); + + // Launch initialized CUTLASS kernel + status = gemm_op(); + CUTLASS_CHECK(status); + + // Create instantiation for device reference gemm kernel + cutlass::reference::device::Gemm + gemm_device; + + // Launch device reference gemm kernel + gemm_device(problem_size, + alpha, + tensor_a.device_ref(), + tensor_b.device_ref(), + beta, + tensor_c.device_ref(), + tensor_ref_d.device_ref()); + + // Wait for kernels to finish + cudaDeviceSynchronize(); + + // Copy output data from CUTLASS and reference kernel to host for comparison + tensor_d.sync_host(); + tensor_ref_d.sync_host(); + + // Check if output from CUTLASS kernel and reference kernel are equal or not + bool passed = cutlass::reference::host::TensorEquals( + tensor_d.host_view(), + tensor_ref_d.host_view()); + + std::cout << (passed ? "Passed" : "Failed") << std::endl; + + return (passed ? 0 : -1); +} + +int main() { + // Turing Tensor Core operations exposed with mma.sync and ldmatrix are first available + // in CUDA 10.2. + // + // CUTLASS must be compiled with CUDA 10.2 Toolkit to run these examples. + if (!(__CUDACC_VER_MAJOR__ > 10 || (__CUDACC_VER_MAJOR__ == 10 && __CUDACC_VER_MINOR__ >= 2))) { + std::cerr << "Turing Tensor Core operations must be compiled with CUDA 10.2 Toolkit or later." << std::endl; + + // Returning zero so this test passes when built on older Toolkits. + return 0; + } + else { + return run(); + } +} +