ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 02:02:46 +00:00
Code Issues Packages Projects Releases Wiki Activity

rename

Browse Source 
[ROCm/composable_kernel commit: 6790b8f3cc]
This commit is contained in:
Chao Liu
2018-11-20 10:43:37 -06:00
parent 3ac454af51
commit cb8367c0eb
7 changed files with 24 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
driver/conv.cu
View File

@@ -5,7 +5,8 @@
#include "nvToolsExt.h"
#include "tensor.hpp"
#include "constant_tensor_descriptor.cuh"
#include "device_direct_convolution_3.cuh"
#include "device_direct_convolution_1.cuh"
#include "device_direct_convolution_2.cuh"
template <class T>
struct GeneratorConstant
@@ -177,7 +178,7 @@ int main()
for(int i = 0; i < 20; ++i)
{
device_convolution(in_desc, in, wei_desc, wei, out_desc, out_device);
device_direct_convolution_2(in_desc, in, wei_desc, wei, out_desc, out_device);
}
#if 0

17
driver/device_direct_convolution_3.cuh → driver/device_direct_convolution_1.cuh
View File

@@ -1,8 +1,8 @@
#pragma once
#include "direct_convolution_3.cuh"
#include "gridwise_direct_convolution_1.cuh"
template <class T, class InDesc, class WeiDesc, class OutDesc>
void device_convolution(
void device_direct_convolution_1(
InDesc, const Tensor<T>& in, WeiDesc, const Tensor<T>& wei, OutDesc, Tensor<T>& out)
{
std::size_t data_sz = sizeof(T);
@@ -26,16 +26,12 @@ void device_convolution(
constexpr auto out_desc = OutDesc{};
constexpr unsigned OutTileSizeH = 2;
constexpr unsigned OutTileSizeW = 2;
constexpr unsigned NPerBlock = 2;
constexpr unsigned KPerBlock = 32;
constexpr unsigned NPerBlock = 1;
constexpr unsigned KPerBlock = 4;
constexpr unsigned CPerBlock = 2;
constexpr unsigned YPerBlock = 1;
constexpr unsigned YPerBlock = 8;
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;
@@ -70,9 +66,6 @@ void device_convolution(
CPerBlock,
YPerBlock,
XPerBlock,
NPerThread,
KPerThread,
CPerThread,
NBlockOpLen0,
NBlockOpLen1,
NBlockOpLen2,

17
driver/device_direct_convolution_2.cuh
View File

@@ -1,8 +1,8 @@
#pragma once
#include "direct_convolution_2.cuh"
#include "gridwise_direct_convolution_2.cuh"
template <class T, class InDesc, class WeiDesc, class OutDesc>
void device_convolution(
void device_direct_convolution_2(
InDesc, const Tensor<T>& in, WeiDesc, const Tensor<T>& wei, OutDesc, Tensor<T>& out)
{
std::size_t data_sz = sizeof(T);
@@ -26,12 +26,16 @@ void device_convolution(
constexpr auto out_desc = OutDesc{};
constexpr unsigned OutTileSizeH = 2;
constexpr unsigned OutTileSizeW = 2;
constexpr unsigned NPerBlock = 1;
constexpr unsigned KPerBlock = 4;
constexpr unsigned NPerBlock = 2;
constexpr unsigned KPerBlock = 32;
constexpr unsigned CPerBlock = 2;
constexpr unsigned YPerBlock = 8;
constexpr unsigned YPerBlock = 1;
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;
@@ -66,6 +70,9 @@ void device_convolution(
CPerBlock,
YPerBlock,
XPerBlock,
NPerThread,
KPerThread,
CPerThread,
NBlockOpLen0,
NBlockOpLen1,
NBlockOpLen2,

2
src/include/blockwise_convolution.cuh → src/include/blockwise_direct_convolution.cuh
View File

@@ -1,7 +1,7 @@
#pragma once
#include "constant_tensor_descriptor.cuh"
#include "threadwise_tensor_op.cuh"
#include "threadwise_convolution.cuh"
#include "threadwise_direct_convolution.cuh"
template <class TFloat,
class InBlockDesc,

2
src/include/direct_convolution_2.cuh → src/include/gridwise_direct_convolution_1.cuh
View File

@@ -1,7 +1,7 @@
#pragma once
#include "constant_tensor_descriptor.cuh"
#include "blockwise_tensor_op.cuh"
#include "blockwise_convolution.cuh"
#include "blockwise_direct_convolution.cuh"
template <class TFloat,
class InGlobalDesc,

4
src/include/direct_convolution_3.cuh → src/include/gridwise_direct_convolution_2.cuh
View File

@@ -1,9 +1,9 @@
#pragma once
#include "constant_tensor_descriptor.cuh"
#include "blockwise_tensor_op.cuh"
#include "blockwise_convolution.cuh"
#include "blockwise_direct_convolution.cuh"
#include "threadwise_tensor_op.cuh"
#include "threadwise_convolution.cuh"
#include "threadwise_direct_convolution.cuh"
template <class TFloat,
class InGlobalDesc,

0
src/include/threadwise_convolution.cuh → src/include/threadwise_direct_convolution.cuh
View File