Skip to content

Commit

Permalink
Rounding Convert and ConvertSat added. (NVIDIA#1068)
Browse files Browse the repository at this point in the history 
* Rounding Convert and ConvertSat added.
* Add a Converter helper class to avoid excessive number of `enable_if`s.
* Add documentation for `Convert[Sat][Norm]` function family.
* Add ./include to directory list processed by Doxygen.

Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>
  • Loading branch information
@mzient @haoxintong
mzient authored and haoxintong committed Jul 16, 2019
1 parent dc3a991 commit 1490468
Show file tree
Hide file tree
Showing 4 changed files with 350 additions and 77 deletions.
2 changes: 1 addition & 1 deletion Doxyfile
View file
Expand Up @@ -771,7 +771,7 @@ WARN_LOGFILE =
# spaces. See also FILE_PATTERNS and EXTENSION_MAPPING
# Note: If this tag is empty the current directory is searched.

INPUT = ./dali ./docs
INPUT = ./dali ./include ./docs

# This tag can be used to specify the character encoding of the source files
# that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses
Expand Down
67 changes: 46 additions & 21 deletions dali/core/convert_cuda_test.cu
View file
Expand Up @@ -14,38 +14,63 @@

#include <gtest/gtest.h>
#include "dali/core/convert.h"
#include "dali/core/math_util.h"
#include "dali/core/convert_test_static.h"
#include "dali/test/device_test.h"

namespace dali {

TEST(ConvertNormCUDA, float2int) {
EXPECT_EQ(ConvertNorm<uint8_t>(0.0f), 0);
EXPECT_EQ(ConvertNorm<uint8_t>(0.499f), 127);
EXPECT_EQ(ConvertNorm<uint8_t>(1.0f), 255);
EXPECT_EQ(ConvertNorm<int8_t>(1.0f), 127);
EXPECT_EQ(ConvertNorm<int8_t>(0.499f), 63);
EXPECT_EQ(ConvertNorm<int8_t>(-1.0f), -127);
DEVICE_TEST(ConvertSatCUDA_Dev, float2int, 110, 1024) {
float f = ldexpf(threadIdx.x-512.0f, blockIdx.x-10);
float integral;
float fract = modff(f, &integral);
if (fract == 0.5f || fract == -0.5f)
return;
double rounded = roundf(f);
int64_t clamped = clamp<double>(rounded, -128, 127);
DEV_EXPECT_EQ(ConvertSat<int8_t>(f), clamped);
clamped = clamp<double>(rounded, 0, 255);
DEV_EXPECT_EQ(ConvertSat<uint8_t>(f), clamped);
clamped = clamp<double>(rounded, -0x8000, 0x7fff);
DEV_EXPECT_EQ(ConvertSat<int16_t>(f), clamped);
clamped = clamp<double>(rounded, 0, 0xffff);
DEV_EXPECT_EQ(ConvertSat<uint16_t>(f), clamped);
clamped = clamp<double>(rounded, int32_t(~0x7fffffff), 0x7fffffff);
DEV_EXPECT_EQ(ConvertSat<int32_t>(f), clamped);
clamped = clamp<double>(rounded, 0, 0xffffffffu);
DEV_EXPECT_EQ(ConvertSat<uint32_t>(f), clamped);
}

DEVICE_TEST(ConvertNormCUDA_Dev, float2int, 1, 1) {
DEV_EXPECT_EQ(ConvertNorm<uint8_t>(0.0f), 0);
DEV_EXPECT_EQ(ConvertNorm<uint8_t>(0.499f), 127);
DEV_EXPECT_EQ(ConvertNorm<uint8_t>(1.0f), 255);
DEV_EXPECT_EQ(ConvertNorm<int8_t>(1.0f), 127);
DEV_EXPECT_EQ(ConvertNorm<int8_t>(0.499f), 63);
DEV_EXPECT_EQ(ConvertNorm<int8_t>(-1.0f), -127);


EXPECT_EQ(ConvertNorm<uint16_t>(0.0f), 0);
EXPECT_EQ(ConvertNorm<uint16_t>(1.0f), 0xffff);
EXPECT_EQ(ConvertNorm<int16_t>(1.0f), 0x7fff);
EXPECT_EQ(ConvertNorm<int16_t>(-1.0f), -0x7fff);
DEV_EXPECT_EQ(ConvertNorm<uint16_t>(0.0f), 0);
DEV_EXPECT_EQ(ConvertNorm<uint16_t>(1.0f), 0xffff);
DEV_EXPECT_EQ(ConvertNorm<int16_t>(1.0f), 0x7fff);
DEV_EXPECT_EQ(ConvertNorm<int16_t>(-1.0f), -0x7fff);
}

TEST(ConvertSatNormCUDA, float2int) {
EXPECT_EQ(ConvertSatNorm<uint8_t>(2.0f), 255);
EXPECT_EQ(ConvertSatNorm<uint8_t>(0.499f), 127);
EXPECT_EQ(ConvertSatNorm<uint8_t>(-2.0f), 0);
EXPECT_EQ(ConvertSatNorm<int8_t>(2.0f), 127);
EXPECT_EQ(ConvertSatNorm<int8_t>(0.499f), 63);
EXPECT_EQ(ConvertSatNorm<int8_t>(-2.0f), -128);
DEVICE_TEST(ConvertSatNorm_Dev, float2int, 1, 1) {
DEV_EXPECT_EQ(ConvertSatNorm<uint8_t>(2.0f), 255);
DEV_EXPECT_EQ(ConvertSatNorm<uint8_t>(0.499f), 127);
DEV_EXPECT_EQ(ConvertSatNorm<uint8_t>(-2.0f), 0);
DEV_EXPECT_EQ(ConvertSatNorm<int8_t>(2.0f), 127);
DEV_EXPECT_EQ(ConvertSatNorm<int8_t>(0.499f), 63);
DEV_EXPECT_EQ(ConvertSatNorm<int8_t>(-2.0f), -128);
DEV_EXPECT_EQ(ConvertSatNorm<uint8_t>(0.4f/255), 0);
DEV_EXPECT_EQ(ConvertSatNorm<uint8_t>(0.6f/255), 1);

EXPECT_EQ(ConvertSatNorm<int16_t>(2.0f), 0x7fff);
EXPECT_EQ(ConvertSatNorm<int16_t>(-2.0f), -0x8000);
DEV_EXPECT_EQ(ConvertSatNorm<int16_t>(2.0f), 0x7fff);
DEV_EXPECT_EQ(ConvertSatNorm<int16_t>(-2.0f), -0x8000);
}

TEST(ConvertNormCUDA, int2float) {
TEST(ConvertNorm_CUDA_Host, int2float) {
EXPECT_EQ((ConvertNorm<float, uint8_t>(255)), 1.0f);
EXPECT_NEAR((ConvertNorm<float, uint8_t>(127)), 1.0f*127/255, 1e-7f);
EXPECT_EQ((ConvertNorm<float, int8_t>(127)), 1.0f);
Expand Down
34 changes: 34 additions & 0 deletions dali/core/convert_test.cc
View file
Expand Up @@ -15,9 +15,41 @@
#include <gtest/gtest.h>
#include "dali/core/convert.h"
#include "dali/core/convert_test_static.h"
#include "dali/core/math_util.h"

namespace dali {


TEST(ConvertSat, float2int) {
for (int exp = -10; exp < 100; exp++) {
for (float sig = -256; sig <= 256; sig++) {
float f = ldexpf(sig, exp);
float integral;
float fract = modff(f, &integral);
if (fract == 0.5f || fract == -0.5f)
continue;
double rounded = roundf(f);
int64_t clamped = clamp<double>(rounded, -128, 127);
ASSERT_EQ(ConvertSat<int8_t>(f), clamped) << " with f = " << f;
clamped = clamp<double>(rounded, 0, 255);
ASSERT_EQ(ConvertSat<uint8_t>(f), clamped) << " with f = " << f;
clamped = clamp<double>(rounded, -0x8000, 0x7fff);
ASSERT_EQ(ConvertSat<int16_t>(f), clamped) << " with f = " << f;
clamped = clamp<double>(rounded, 0, 0xffff);
ASSERT_EQ(ConvertSat<uint16_t>(f), clamped) << " with f = " << f;
clamped = clamp<double>(rounded, int32_t(~0x7fffffff), 0x7fffffff);
ASSERT_EQ(ConvertSat<int32_t>(f), clamped) << " with f = " << f;
clamped = clamp<double>(rounded, 0, 0xffffffffu);
ASSERT_EQ(ConvertSat<uint32_t>(f), clamped) << " with f = " << f;
}
}
}

TEST(ConvertNorm, int2int) {
EXPECT_EQ((ConvertNorm<uint8_t, uint8_t>(0)), 0);
EXPECT_EQ((ConvertNorm<uint8_t, int8_t>(127)), 255);
}

TEST(ConvertNorm, float2int) {
EXPECT_EQ(ConvertNorm<uint8_t>(0.0f), 0);
EXPECT_EQ(ConvertNorm<uint8_t>(0.499f), 127);
Expand All @@ -40,6 +72,8 @@ TEST(ConvertSatNorm, float2int) {
EXPECT_EQ(ConvertSatNorm<int8_t>(2.0f), 127);
EXPECT_EQ(ConvertSatNorm<int8_t>(0.499f), 63);
EXPECT_EQ(ConvertSatNorm<int8_t>(-2.0f), -128);
EXPECT_EQ(ConvertSatNorm<uint8_t>(0.4f/255), 0);
EXPECT_EQ(ConvertSatNorm<uint8_t>(0.6f/255), 1);

EXPECT_EQ(ConvertSatNorm<int16_t>(2.0f), 0x7fff);
EXPECT_EQ(ConvertSatNorm<int16_t>(-2.0f), -0x8000);
Expand Down

0 comments on commit 1490468

Please sign in to comment.