Skip to content
This repository has been archived by the owner on Nov 17, 2023. It is now read-only.

[MXNET-380] count_include_pad argument for Avg Pooling #11021

Merged
merged 6 commits into from
Jun 18, 2018
Merged

[MXNET-380] count_include_pad argument for Avg Pooling #11021

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
b3fdf26
add count_include_pad argument
May 22, 2018
5b5b702
add cound_include_pad in cudnn pooling and corresponding tests
Jun 1, 2018
dec157a
add gluon support for the new option
Jun 2, 2018
7935e8e
switch to built-in functions for artificial NaN
Jun 4, 2018
db6279d
add doc for gluon
Jun 9, 2018
5a07a43
change isAvg/getAvg to is_avg/get_avg
Jun 11, 2018
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
add cound_include_pad in cudnn pooling and corresponding tests
  • Loading branch information
Hao Jin committed Jun 11, 2018
commit 5b5b7028c633a68432877d84439f065ca53c87f2
59 changes: 0 additions & 59 deletions scala-package/macros/src/main/scala/org/apache/mxnet/SymbolMacro.scala
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -178,65 +178,6 @@ private[mxnet] object SymbolImplMacros {
result
}

// Convert C++ Types to Scala Types
def typeConversion(in : String, argType : String = "") : String = {
in match {
case "Shape(tuple)" | "ShapeorNone" => "org.apache.mxnet.Shape"
case "Symbol" | "NDArray" | "NDArray-or-Symbol" => "org.apache.mxnet.Symbol"
case "Symbol[]" | "NDArray[]" | "NDArray-or-Symbol[]" | "SymbolorSymbol[]"
=> "Array[org.apache.mxnet.Symbol]"
case "float" | "real_t" | "floatorNone" => "org.apache.mxnet.Base.MXFloat"
case "int" | "intorNone" | "int(non-negative)" => "Int"
case "long" | "long(non-negative)" => "Long"
case "double" | "doubleorNone" => "Double"
case "string" => "String"
case "boolean" | "booleanorNone" => "Boolean"
case "tupleof<float>" | "tupleof<double>" | "ptr" | "" => "Any"
case default => throw new IllegalArgumentException(
s"Invalid type for args: $default, $argType")
}
}


/**
* By default, the argType come from the C++ API is a description more than a single word
* For Example:
* <C++ Type>, <Required/Optional>, <Default=>
* The three field shown above do not usually come at the same time
* This function used the above format to determine if the argument is
* optional, what is it Scala type and possibly pass in a default value
* @param argType Raw arguement Type description
* @return (Scala_Type, isOptional)
*/
def argumentCleaner(argType : String) : (String, Boolean) = {
val spaceRemoved = argType.replaceAll("\\s+", "")
var commaRemoved : Array[String] = new Array[String](0)
// Deal with the case e.g: stype : {'csr', 'default', 'row_sparse'}
if (spaceRemoved.charAt(0)== '{') {
val endIdx = spaceRemoved.indexOf('}')
commaRemoved = spaceRemoved.substring(endIdx + 1).split(",")
commaRemoved(0) = "string"
} else {
commaRemoved = spaceRemoved.split(",")
}
// Optional Field
if (commaRemoved.length >= 3) {
// arg: Type, optional, default = Null
require(commaRemoved(1).equals("optional"))
require(commaRemoved(2).startsWith("default="))
(typeConversion(commaRemoved(0), argType), true)
} else if (commaRemoved.length == 2 || commaRemoved.length == 1) {
val tempType = typeConversion(commaRemoved(0), argType)
val tempOptional = tempType.equals("org.apache.mxnet.Symbol")
(tempType, tempOptional)
} else {
throw new IllegalArgumentException(
s"Unrecognized arg field: $argType, ${commaRemoved.length}")
}

}


// List and add all the atomic symbol functions to current module.
private def initSymbolModule(): List[SymbolFunction] = {
val opNames = ListBuffer.empty[String]
Expand Down
8 changes: 6 additions & 2 deletions src/operator/nn/cudnn/cudnn_pooling-inl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,7 +51,11 @@ class CuDNNPoolingOp {
mode_ = CUDNN_POOLING_MAX;
break;
case pool_enum::kAvgPooling:
mode_ = CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
if (param_.count_include_pad.has_value() && !param_.count_include_pad.value()) {
mode_ = CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING;
} else {
mode_ = CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
}
break;
default:
LOG(FATAL) << "Not implmented";
Expand Down Expand Up @@ -263,7 +267,7 @@ class CuDNNPoolingOp {
&(pad_vec[0]),
&(stride_vec[0])));
#else
LOG(FATAL) << "3D pooling only support CUDNN v5 and abouve";
LOG(FATAL) << "3D pooling only support CUDNN v5 and above";
#endif
}
}
Expand Down
4 changes: 2 additions & 2 deletions src/operator/nn/pool.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -266,7 +266,7 @@ __global__ void pool_sum_2d_gpu_kernel(const int nthreads, const DType* in_data,
if (getAvg && !count_include_pad) {
pool_size = (hend - hstart) * (wend - wstart);
}
DType sum = 0;
DType sum = 0.0f;
const DType* out_slice = in_data + (n * channels + c) * height * width;
for (int h = hstart; h < hend; ++h) {
for (int w = wstart; w < wend; ++w) {
Expand Down Expand Up @@ -322,7 +322,7 @@ __global__ void pool_sum_3d_gpu_kernel(const int nthreads, const DType* in_data,
}
}
}
out_data[index] = a_root_p<DType, p>::Map(sum);
out_data[index] = (pool_size == 0) ? DType(0.0f / 0) : a_root_p<DType, p>::Map(sum);
}
}

Expand Down
4 changes: 3 additions & 1 deletion src/operator/nn/pool.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -340,7 +340,9 @@ inline void pool_sum_3d_cpu(const DType* in_data, const TShape& ishape, const TS
}
}
}
out_data[(pd*pooled_height+ph)*pooled_width+pw] = a_root_p<DType, p>::Map(sum);
out_data[(pd*pooled_height+ph)*pooled_width+pw] = (pool_size == 0) ?
DType(0.0f / 0) :
a_root_p<DType, p>::Map(sum);
}
}
}
Expand Down
90 changes: 57 additions & 33 deletions tests/python/gpu/test_operator_gpu.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -740,8 +740,8 @@ def test_pooling_with_type():

@with_seed()
def test_pooling_versions():
def test_pooling_versions_helper(pool_op_list, data, kernel, pool_type, pad, stride,
pooling_convention='valid', global_pool=False, p_value=2):
def test_pooling_versions_helper(pool_op_list, data, kernel, pool_type, pad, stride, pooling_convention='valid',
global_pool=False, p_value=2, count_include_pad=True, tol=None):
ctx_list = []
sym_list = []
# PoolingV1 cpu
Expand All @@ -765,61 +765,69 @@ def test_pooling_versions_helper(pool_op_list, data, kernel, pool_type, pad, str
ctx_list.append({'ctx': mx.cpu(0), 'pool_data': data, 'type_dict': {'pool_data': np.float32}})
if not global_pool:
sym_list.append(mx.sym.Pooling(kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention=pooling_convention, name='pool', p_value=p_value))
pooling_convention=pooling_convention, name='pool',
p_value=p_value, count_include_pad=count_include_pad))
else:
sym_list.append(mx.sym.Pooling(kernel=kernel, pool_type=pool_type, global_pool=True, name='pool', p_value=p_value))
sym_list.append(mx.sym.Pooling(kernel=kernel, pool_type=pool_type, global_pool=True, name='pool',
p_value=p_value, count_include_pad=count_include_pad))
# Pooling gpu
if 'pool_gpu' in pool_op_list:
ctx_list.append({'ctx': mx.gpu(0), 'pool_data': data, 'type_dict': {'pool_data': np.float32}})
if not global_pool:
sym_list.append(mx.sym.Pooling(kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention=pooling_convention, cudnn_off=True, name='pool', p_value=p_value))
pooling_convention=pooling_convention, cudnn_off=True, name='pool',
p_value=p_value, count_include_pad=count_include_pad))
else:
sym_list.append(mx.sym.Pooling(kernel=kernel, pool_type=pool_type, global_pool=True, cudnn_off=True,
name='pool', p_value=p_value))
name='pool', p_value=p_value, count_include_pad=count_include_pad))
# CuDNNPooling
if 'pool_cudnn' in pool_op_list:
ctx_list.append({'ctx': mx.gpu(0), 'pool_data': data, 'type_dict': {'pool_data': np.float32}})
if not global_pool:
sym_list.append(mx.sym.Pooling(kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention=pooling_convention, p_value=p_value, cudnn_off=False, name='pool'))
pooling_convention=pooling_convention, p_value=p_value, cudnn_off=False,
name='pool', count_include_pad=count_include_pad))
else:
sym_list.append(mx.sym.Pooling(kernel=kernel, pool_type=pool_type, global_pool=True, p_value=p_value,
cudnn_off=False, name='pool'))
check_consistency(sym_list, ctx_list)
cudnn_off=False, name='pool', count_include_pad=count_include_pad))
check_consistency(sym_list, ctx_list, equal_nan=(not count_include_pad), tol=tol)

def test_1d_pooling(pool_type, p_value=2):
def test_1d_pooling(pool_type, p_value=2, count_include_pad=True):
data = (2, 3, 20)
kernel = (4,)
pad = (0,)
stride = (1,)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='valid', global_pool=False, p_value=p_value)
pooling_convention='valid', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (2,)
stride = (2,)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='valid', global_pool=False, p_value=p_value)
pooling_convention='valid', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (0,)
stride = (1,)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False, p_value=p_value)
pooling_convention='full', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (2,)
stride = (2,)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False, p_value=p_value)
pooling_convention='full', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
global_pool=True, p_value=p_value)
global_pool=True, p_value=p_value, count_include_pad=count_include_pad)

def test_2d_pooling(pool_type, p_value=2):
def test_2d_pooling(pool_type, p_value=2, count_include_pad=True):
data = (2, 3, 20, 20)
kernel = (4, 5)
pad = (0, 0)
Expand All @@ -831,14 +839,15 @@ def test_2d_pooling(pool_type, p_value=2):
else:
test_pooling_versions_helper(pool_op_list=['pool_v1_cpu', 'pool_v1_gpu', 'pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='valid', global_pool=False)
pooling_convention='valid', global_pool=False, count_include_pad=count_include_pad)

# pool_v1 has bugs when pad is not 0, do not test PoolingV1 here
pad = (2, 3)
stride = (2, 3)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='valid', global_pool=False, p_value=p_value)
pooling_convention='valid', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (0, 0)
stride = (1, 1)
Expand All @@ -847,16 +856,24 @@ def test_2d_pooling(pool_type, p_value=2):
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False, p_value=p_value)
else:
test_pooling_versions_helper(pool_op_list=['pool_v1_cpu', 'pool_v1_gpu', 'pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False)
if count_include_pad:
test_pooling_versions_helper(pool_op_list=['pool_v1_cpu', 'pool_v1_gpu', 'pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False,
count_include_pad=count_include_pad)
else:
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False,
count_include_pad=count_include_pad)

# pool_v1 has bugs when pad is not 0, do not test PoolingV1 here
pad = (2, 3)
stride = (2, 3)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False, p_value=p_value)
pooling_convention='full', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

if pool_type == 'lp':
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
Expand All @@ -865,55 +882,62 @@ def test_2d_pooling(pool_type, p_value=2):
else:
test_pooling_versions_helper(pool_op_list=['pool_v1_cpu', 'pool_v1_gpu', 'pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
global_pool=True)
global_pool=True, count_include_pad=count_include_pad)

def test_3d_pooling(pool_type, p_value=2):
data = (2, 3, 20, 20, 20)
def test_3d_pooling(pool_type, p_value=2, count_include_pad=True):
data = (1, 1, 20, 20, 20)
kernel = (4, 5, 3)
pad = (0, 0, 0)
stride = (1, 1, 1)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='valid', global_pool=False, p_value=p_value)
pooling_convention='valid', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (2, 3, 3)
stride = (2, 3, 1)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='valid', global_pool=False, p_value=p_value)
pooling_convention='valid', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (0, 0, 0)
stride = (1, 1, 1)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False, p_value=p_value)
pooling_convention='full', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

pad = (2, 3, 3)
stride = (2, 3, 1)
test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
pooling_convention='full', global_pool=False, p_value=p_value)
pooling_convention='full', global_pool=False, p_value=p_value,
count_include_pad=count_include_pad)

test_pooling_versions_helper(pool_op_list=['pool_cpu', 'pool_gpu', 'pool_cudnn'],
data=data, kernel=kernel, pad=pad, stride=stride, pool_type=pool_type,
global_pool=True, p_value=p_value)
global_pool=True, p_value=p_value, count_include_pad=count_include_pad)

test_1d_pooling('max')
test_1d_pooling('avg')
test_1d_pooling('avg', count_include_pad=True)
test_1d_pooling('avg', count_include_pad=False)
test_1d_pooling('sum')
test_1d_pooling('lp', p_value=1)
test_1d_pooling('lp', p_value=2)
test_1d_pooling('lp', p_value=3)

test_2d_pooling('max')
test_2d_pooling('avg')
test_2d_pooling('avg', count_include_pad=True)
test_2d_pooling('avg', count_include_pad=False)
test_2d_pooling('sum')
test_2d_pooling('lp', p_value=1)
test_2d_pooling('lp', p_value=2)
test_2d_pooling('lp', p_value=3)

test_3d_pooling('max')
test_3d_pooling('avg')
test_3d_pooling('avg', count_include_pad=True)
test_3d_pooling('avg', count_include_pad=False)
test_3d_pooling('sum')
test_3d_pooling('lp', p_value=1)
test_3d_pooling('lp', p_value=2)
Expand Down