adding the cutlass examples

1 files changed, 152 insertions, 0 deletions

diff --git a/cutlass-example/gemm.h b/cutlass-example/gemm.h
new file mode 100644
index 0000000..18dfdf6
--- /dev/null
+++ b/cutlass-example/gemm.h
@@ -0,0 +1,152 @@
+/***************************************************************************************************
+* 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);
+}
+
+////////////////////////////////////////////////////////////////////////////////////////////////////