From 69a23647d8fe8d4a2a853a6b7a55c63283497bdf Mon Sep 17 00:00:00 2001
From: Chao Liu <lc.roy86@gmail.com>
Date: Tue, 20 Nov 2018 10:51:28 -0600
Subject: [PATCH] refactor

[ROCm/composable_kernel commit: a21b0d27a52794fcf23e34f496559c0f1adc1a0d]
---
 driver/conv.cu                                |  2 +-
 driver/device_direct_convolution_1.cuh        | 34 ++++++++--------
 driver/device_direct_convolution_2.cuh        | 40 +++++++++----------
 src/include/gridwise_direct_convolution_1.cuh | 12 +++---
 src/include/gridwise_direct_convolution_2.cuh | 12 +++---
 5 files changed, 50 insertions(+), 50 deletions(-)

diff --git a/driver/conv.cu b/driver/conv.cu
index 966f94dfc7..b1d9308b97 100644
--- a/driver/conv.cu
+++ b/driver/conv.cu
@@ -178,7 +178,7 @@ int main()
 
     for(int i = 0; i < 20; ++i)
     {
-        device_direct_convolution_2(in_desc, in, wei_desc, wei, out_desc, out_device);
+        device_direct_convolution_1(in_desc, in, wei_desc, wei, out_desc, out_device);
     }
 
 #if 0
diff --git a/driver/device_direct_convolution_1.cuh b/driver/device_direct_convolution_1.cuh
index 7a0fc29fbe..5527dd3946 100644
--- a/driver/device_direct_convolution_1.cuh
+++ b/driver/device_direct_convolution_1.cuh
@@ -55,23 +55,23 @@ void device_direct_convolution_1(
     cudaEventCreate(&start);
     cudaEventRecord(start, 0);
 
-    gridwise_convolution<T,
-                         InDesc,
-                         WeiDesc,
-                         OutDesc,
-                         OutTileSizeH,
-                         OutTileSizeW,
-                         NPerBlock,
-                         KPerBlock,
-                         CPerBlock,
-                         YPerBlock,
-                         XPerBlock,
-                         NBlockOpLen0,
-                         NBlockOpLen1,
-                         NBlockOpLen2,
-                         NBlockOpLen3,
-                         BlockSize,
-                         GridSize>
+    gridwise_direct_convolution_1<T,
+                                  InDesc,
+                                  WeiDesc,
+                                  OutDesc,
+                                  OutTileSizeH,
+                                  OutTileSizeW,
+                                  NPerBlock,
+                                  KPerBlock,
+                                  CPerBlock,
+                                  YPerBlock,
+                                  XPerBlock,
+                                  NBlockOpLen0,
+                                  NBlockOpLen1,
+                                  NBlockOpLen2,
+                                  NBlockOpLen3,
+                                  BlockSize,
+                                  GridSize>
         <<<grid_dim, block_dim>>>(InDesc{},
                                   static_cast<T*>(in_device_buf.GetDeviceBuffer()),
                                   WeiDesc{},
diff --git a/driver/device_direct_convolution_2.cuh b/driver/device_direct_convolution_2.cuh
index 12c27b36cb..4998149716 100644
--- a/driver/device_direct_convolution_2.cuh
+++ b/driver/device_direct_convolution_2.cuh
@@ -59,26 +59,26 @@ void device_direct_convolution_2(
     cudaEventCreate(&start);
     cudaEventRecord(start, 0);
 
-    gridwise_convolution<T,
-                         InDesc,
-                         WeiDesc,
-                         OutDesc,
-                         OutTileSizeH,
-                         OutTileSizeW,
-                         NPerBlock,
-                         KPerBlock,
-                         CPerBlock,
-                         YPerBlock,
-                         XPerBlock,
-                         NPerThread,
-                         KPerThread,
-                         CPerThread,
-                         NBlockOpLen0,
-                         NBlockOpLen1,
-                         NBlockOpLen2,
-                         NBlockOpLen3,
-                         BlockSize,
-                         GridSize>
+    gridwise_direct_convolution_2<T,
+                                  InDesc,
+                                  WeiDesc,
+                                  OutDesc,
+                                  OutTileSizeH,
+                                  OutTileSizeW,
+                                  NPerBlock,
+                                  KPerBlock,
+                                  CPerBlock,
+                                  YPerBlock,
+                                  XPerBlock,
+                                  NPerThread,
+                                  KPerThread,
+                                  CPerThread,
+                                  NBlockOpLen0,
+                                  NBlockOpLen1,
+                                  NBlockOpLen2,
+                                  NBlockOpLen3,
+                                  BlockSize,
+                                  GridSize>
         <<<grid_dim, block_dim>>>(InDesc{},
                                   static_cast<T*>(in_device_buf.GetDeviceBuffer()),
                                   WeiDesc{},
diff --git a/src/include/gridwise_direct_convolution_1.cuh b/src/include/gridwise_direct_convolution_1.cuh
index de3b31bc0a..52949b49c1 100644
--- a/src/include/gridwise_direct_convolution_1.cuh
+++ b/src/include/gridwise_direct_convolution_1.cuh
@@ -20,12 +20,12 @@ template <class TFloat,
           unsigned NBlockOpLen3,
           unsigned BlockSize,
           unsigned GridSize>
-__global__ void gridwise_convolution(InGlobalDesc,
-                                     TFloat* const __restrict__ p_in_global,
-                                     WeiGlobalDesc,
-                                     TFloat* const __restrict__ p_wei_global,
-                                     OutGlobalDesc,
-                                     TFloat* __restrict__ p_out_global)
+__global__ void gridwise_direct_convolution_1(InGlobalDesc,
+                                              TFloat* const __restrict__ p_in_global,
+                                              WeiGlobalDesc,
+                                              TFloat* const __restrict__ p_wei_global,
+                                              OutGlobalDesc,
+                                              TFloat* __restrict__ p_out_global)
 {
     constexpr auto I0 = Index<0>{};
     constexpr auto I1 = Index<1>{};
diff --git a/src/include/gridwise_direct_convolution_2.cuh b/src/include/gridwise_direct_convolution_2.cuh
index 135632dca6..e4da59b1fe 100644
--- a/src/include/gridwise_direct_convolution_2.cuh
+++ b/src/include/gridwise_direct_convolution_2.cuh
@@ -25,12 +25,12 @@ template <class TFloat,
           unsigned NBlockOpLen3,
           unsigned BlockSize,
           unsigned GridSize>
-__global__ void gridwise_convolution(InGlobalDesc,
-                                     TFloat* const __restrict__ p_in_global,
-                                     WeiGlobalDesc,
-                                     TFloat* const __restrict__ p_wei_global,
-                                     OutGlobalDesc,
-                                     TFloat* __restrict__ p_out_global)
+__global__ void gridwise_direct_convolution_2(InGlobalDesc,
+                                              TFloat* const __restrict__ p_in_global,
+                                              WeiGlobalDesc,
+                                              TFloat* const __restrict__ p_wei_global,
+                                              OutGlobalDesc,
+                                              TFloat* __restrict__ p_out_global)
 {
     constexpr auto I0 = Index<0>{};
     constexpr auto I1 = Index<1>{};