Skip to content

unrepresentable indices in max_pooling/max_unpooling in fp16 target leads to illegal memory accesses #15996

@YashasSamaga

Description

@YashasSamaga
System information (version)
  • OpenCV => master @ ad0ab41
  • Operating System / Platform => Ubuntu 18.04 64 Bit
  • Compiler => GCC 7.4.0
Detailed description

Half-precision floats are not capable of accurately storing the indices in max_pooling layers for ENet.

max_idx as index_type: 65521
static_cast<__half>(max_idx): inf
static_cast<index_type>(static_cast<__half>(max_idx)): 2147483647

The max_pooling kernel computes the indices as integers but finally stores the index in __half format. The index which is retrieved after casting back to index_type in max_unpooling kernel can potentially leads to illegal memory access.

This renders the FP16 target unusable in networks which required max_unpooling for even reasonably small feature maps.

Steps to reproduce

DNNTestNetwork.ENet/1, where GetParam() = CUDA/CUDA_FP16

Metadata

Metadata

Assignees

Type

No type

Projects

No projects

Relationships

None yet

Development

No branches or pull requests

Issue actions