aboutsummaryrefslogtreecommitdiff
path: root/cutlass-example/gemm.h
diff options
context:
space:
mode:
Diffstat (limited to 'cutlass-example/gemm.h')
-rw-r--r--cutlass-example/gemm.h152
1 files changed, 0 insertions, 152 deletions
diff --git a/cutlass-example/gemm.h b/cutlass-example/gemm.h
deleted file mode 100644
index 18dfdf6..0000000
--- a/cutlass-example/gemm.h
+++ /dev/null
@@ -1,152 +0,0 @@
-/***************************************************************************************************
-* Copyright (c) 2018, 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.
-*
-**************************************************************************************************/
-
-#include <cutlass/cutlass.h>
-#include <gemm_testbed.h>
-
-////////////////////////////////////////////////////////////////////////////////////////////////////
-
-template <typename GemmTraits_>
-static void run_gemm(
- int m,
- int n,
- int k,
- int lda,
- int ldb,
- int ldc,
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type alpha =
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type(1),
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type beta =
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type(0)) {
- typedef cutlass::gemm::Gemm<GemmTraits_> Gemm;
- typename Gemm::Params params;
-
- printf("run_gemm-2:m=%d\n",m);
- printf("run_gemm-2:n=%d\n",n);
- printf("run_gemm-2:k=%d\n",k);
- printf("run_gemm-2:lda=%d\n",lda);
- printf("run_gemm-2:ldb=%d\n",ldb);
- printf("run_gemm-2:ldc=%d\n",ldc);
- printf("run_gemm-2:alpha=%.2f\n",alpha);
- printf("run_gemm-2:beta=%.2f\n",beta);
-
- test::GemmTestbed<
- typename test::GemmTestbedTraits<
- typename GemmTraits_::GemmConfig::ScalarA>::host_type, // AType
- typename test::GemmTestbedTraits<
- typename GemmTraits_::GemmConfig::ScalarB>::host_type, // BType
- typename test::GemmTestbedTraits<
- typename GemmTraits_::Epilogue::ScalarC>::host_type, // CType
- typename test::GemmTestbedTraits<
- typename GemmTraits_::Epilogue::Accumulators::Element>::host_type, // Accumulator
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type // Scalar
- >
- testbed(m,
- n,
- k,
- lda,
- ldb,
- ldc,
- cutlass::convert(GemmTraits_::kLayoutA),
- cutlass::convert(GemmTraits_::kLayoutB),
- alpha,
- beta);
-
- testbed.initialize();
-
- // if (testbed.has_cublas_support()) {
- // EXPECT_TRUE(testbed.verify_host_with_cublas());
- // }
-
- params.initialize(testbed.M(),
- testbed.N(),
- testbed.K(),
- testbed.alpha,
- testbed.ptr_A(),
- testbed.lda(),
- testbed.ptr_B(),
- testbed.ldb(),
- testbed.beta,
- testbed.ptr_C_initial(),
- testbed.ldc(),
- testbed.ptr_computed(),
- testbed.ldc());
-
- printf("SIZE_OF_PARAM=%lu\n",sizeof(params));
- void *ptr =&params;
- for(int kk=0;kk<108;kk++){
- printf("KERNELPARAM:%d:%08x\n",kk,*((((int *) ptr)+kk)));
- }
- printf("m=%lu\n",sizeof(params.m));
- printf("n=%lu\n",sizeof(params.n));
- printf("k=%lu\n",sizeof(params.k));
-// printf("alpha=%d\n",sizeof(params.alpha));
-// printf("beta=%d\n",sizeof(params.beta));
-// printf("d_a=%d\n",sizeof(params.d_a));
-// printf("lda=%d\n",sizeof(params.lda));
-// printf("d_b=%d\n",sizeof(params.d_b));
-// printf("ldb=%d\n",sizeof(params.ldb));
-// printf("d_c=%d\n",sizeof(params.d_c));
-// printf("ldc=%d\n",sizeof(params.ldc));
-// printf("d_d=%d\n",sizeof(params.d_d));
-// printf("ldd=%d\n",sizeof(params.ldd));
- Gemm::launch(params);
-
- cudaError_t result = cudaDeviceSynchronize();
- if(result==cudaSuccess){
- printf("Successfully Launched\n");
- }
- int save=1;
- int myval=testbed.verify_with_host(save,save);
- if(myval==1){
- printf("Result Verified\n");
- }
-
-}
-
-////////////////////////////////////////////////////////////////////////////////////////////////////
-
-template <typename GemmTraits_>
-static void run_gemm(
- int m,
- int n,
- int k,
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type alpha =
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type(1),
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type beta =
- typename test::GemmTestbedTraits<typename GemmTraits_::Epilogue::Scalar>::host_type(0)) {
- int lda = GemmTraits_::kLayoutA == cutlass::MatrixLayout::kColumnMajor ? m : k;
- int ldb = GemmTraits_::kLayoutB == cutlass::MatrixLayout::kColumnMajor ? k : n;
- printf("run_gemm-1:m=%d\n",m);
- printf("run_gemm-1:n=%d\n",n);
- printf("run_gemm-1:k=%d\n",k);
- printf("run_gemm-1:alpha=%.2f\n",alpha);
- printf("run_gemm-1:beta=%.2f\n",beta);
- printf("run_gemm-1:lda=%d\n",lda);
- printf("run_gemm-1:ldb=%d\n",ldb);
- run_gemm<GemmTraits_>(m, n, k, lda, ldb, m, alpha, beta);
-}
-
-////////////////////////////////////////////////////////////////////////////////////////////////////
generated by cgit v1.3 (git 2.53.0) at 2026-07-21 15:28:36 +0000