faster: output skip LDS

driver/conv.cu

@@ -5,7 +5,7 @@
 #include "nvToolsExt.h"
 #include "tensor.hpp"
 #include "constant_tensor_descriptor.cuh"
-#include "device_direct_convolution_2.cuh"
+#include "device_direct_convolution_3.cuh"
 
 template <class T>
 struct GeneratorConstant
@@ -27,8 +27,9 @@ struct GeneratorTensor
     {
 #if 1
         return T(std::rand()) / T(RAND_MAX);
-#elif
-
+#elif 1
+        return 1;
+#elif 0
         std::initializer_list<std::size_t> ls = {static_cast<std::size_t>(is)...};
         return std::accumulate(ls.begin(), ls.end(), std::size_t(0));
 #else
@@ -111,11 +112,12 @@ void host_convolution(const Tensor<T>& in, const Tensor<T>& wei, Tensor<T>& out)
 
 int main()
 {
+
 #if 0
     constexpr unsigned N  = 1;
     constexpr unsigned C  = 1;
-    constexpr unsigned HI = 18;
-    constexpr unsigned WI = 18;
+    constexpr unsigned HI = 34;
+    constexpr unsigned WI = 34;
     constexpr unsigned K  = 1;
     constexpr unsigned S  = 3;
     constexpr unsigned R  = 3;
@@ -127,6 +129,14 @@ int main()
     constexpr unsigned K = 64;
     constexpr unsigned S = 3;
     constexpr unsigned R = 3;
+#elif 0
+    constexpr unsigned N = 1;
+    constexpr unsigned C = 1;
+    constexpr unsigned HI = 18;
+    constexpr unsigned WI = 18;
+    constexpr unsigned K = 1;
+    constexpr unsigned S = 3;
+    constexpr unsigned R = 3;
 #elif 0
     constexpr unsigned N = 2;
     constexpr unsigned C = 3;
@@ -156,18 +166,19 @@ int main()
     Tensor<float> in(make_TensorDescriptor(in_desc));
     Tensor<float> wei(make_TensorDescriptor(wei_desc));
     Tensor<float> out_host(make_TensorDescriptor(out_desc));
+    Tensor<float> out_device(make_TensorDescriptor(out_desc));
 
     int num_thread = std::thread::hardware_concurrency();
 
 #if 1
     in.GenerateTensorValue(GeneratorTensor<float>{}, num_thread);
     wei.GenerateTensorValue(GeneratorTensor<float>{}, num_thread);
-    out_host.GenerateTensorValue(GeneratorConstant<float>{0}, num_thread);
 #endif
 
-    Tensor<float> out_device = out_host;
-
-    device_convolution(in_desc, in, wei_desc, wei, out_desc, out_device);
+    for(int i = 0; i < 20; ++i)
+    {
+        device_convolution(in_desc, in, wei_desc, wei, out_desc, out_device);
+    }
 
 #if 1
     host_convolution(in, wei, out_host);
@@ -192,9 +203,9 @@ int main()
 #endif
 
 #if 0
-    LogRange(std::cout << __func__ << "in : ", in.mData, ",") << std::endl;
-    LogRange(std::cout << __func__ << "wei: ", wei.mData, ",") << std::endl;
-    LogRange(std::cout, out_host.mData, ",") << std::endl;
-    LogRange(std::cout, out_device.mData, ",") << std::endl;
+    LogRange(std::cout << "in : ", in.mData, ",") << std::endl;
+    LogRange(std::cout << "wei: ", wei.mData, ",") << std::endl;
+    LogRange(std::cout << "out_host  : ", out_host.mData, ",") << std::endl;
+    LogRange(std::cout << "out_device: ", out_device.mData, ",") << std::endl;
 #endif
 }

driver/device_direct_convolution_3.cuh

@@ -0,0 +1,98 @@
+#pragma once
+#include "direct_convolution_3.cuh"
+
+template <class T, class InDesc, class WeiDesc, class OutDesc>
+void device_convolution(
+    InDesc, const Tensor<T>& in, WeiDesc, const Tensor<T>& wei, OutDesc, Tensor<T>& out)
+{
+    std::size_t data_sz = sizeof(T);
+    DeviceMem in_device_buf(data_sz * in.mDesc.GetElementSpace());
+    DeviceMem wei_device_buf(data_sz * wei.mDesc.GetElementSpace());
+    DeviceMem out_device_buf(data_sz * out.mDesc.GetElementSpace());
+
+    int num_thread = std::thread::hardware_concurrency();
+
+    in_device_buf.ToDevice(in.mData.data());
+    wei_device_buf.ToDevice(wei.mData.data());
+    out_device_buf.ToDevice(out.mData.data());
+
+    constexpr auto I0 = Index<0>{};
+    constexpr auto I1 = Index<1>{};
+    constexpr auto I2 = Index<2>{};
+    constexpr auto I3 = Index<3>{};
+
+    constexpr auto in_desc          = InDesc{};
+    constexpr auto wei_desc         = WeiDesc{};
+    constexpr auto out_desc         = OutDesc{};
+    constexpr unsigned OutTileSizeH = 2;
+    constexpr unsigned OutTileSizeW = 2;
+    constexpr unsigned NPerBlock    = 2;
+    constexpr unsigned KPerBlock    = 8;
+    constexpr unsigned CPerBlock    = 2;
+    constexpr unsigned YPerBlock    = 4;
+    constexpr unsigned XPerBlock    = 16;
+
+    constexpr unsigned NPerThread = 2;
+    constexpr unsigned KPerThread = 4;
+    constexpr unsigned CPerThread = 2;
+
+    constexpr unsigned NBlockOpLen0 = 1;
+    constexpr unsigned NBlockOpLen1 = 1;
+    constexpr unsigned NBlockOpLen2 = 4;
+    constexpr unsigned NBlockOpLen3 = 32;
+
+    constexpr unsigned BlockSize = 128;
+
+    constexpr unsigned GridSize = (out_desc.GetLength(I0) / NPerBlock) *
+                                  (out_desc.GetLength(I1) / KPerBlock) *
+                                  (out_desc.GetLength(I2) / (OutTileSizeH * YPerBlock)) *
+                                  (out_desc.GetLength(I3) / (OutTileSizeW * XPerBlock));
+
+    dim3 block_dim(BlockSize);
+    dim3 grid_dim(GridSize);
+
+    printf("%s: BlockSize %u, GridSize %u \n", __func__, BlockSize, GridSize);
+
+    cudaEvent_t start, stop;
+    float elapsedTime;
+
+    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>
+        <<<grid_dim, block_dim>>>(InDesc{},
+                                  static_cast<T*>(in_device_buf.GetDeviceBuffer()),
+                                  WeiDesc{},
+                                  static_cast<T*>(wei_device_buf.GetDeviceBuffer()),
+                                  OutDesc{},
+                                  static_cast<T*>(out_device_buf.GetDeviceBuffer()));
+
+    cudaEventCreate(&stop);
+    cudaEventRecord(stop, 0);
+    cudaEventSynchronize(stop);
+
+    cudaEventElapsedTime(&elapsedTime, start, stop);
+    printf("Elapsed time : %f ms\n", elapsedTime);
+
+    checkCudaErrors(cudaGetLastError());
+    out_device_buf.FromDevice(out.mData.data());
+}