Fix color operators.

* BrightnessContrast kernel: loop boundaries * HSV, BrightnessContrast: propagate layout * HSV: enforce input layout Signed-off-by: Michal Zientkiewicz <michalz@nvidia.com>
NVIDIA · Dec 9, 2019 · 0ba8b46 · 0ba8b46
1 parent d5b3f25
commit 0ba8b46
Show file tree

Hide file tree

Showing 5 changed files with 20 additions and 16 deletions.
diff --git a/dali/kernels/imgproc/color_manipulation/brightness_contrast_gpu.h b/dali/kernels/imgproc/color_manipulation/brightness_contrast_gpu.h
@@ -128,8 +128,8 @@ BrightnessContrastKernel(const SampleDescriptor<OutputType, InputType, ndims> *s
   const auto *__restrict__ in = sample.in;
   auto *__restrict__ out = sample.out;
 
-  for (int y = threadIdx.y + block.start.y; y < threadIdx.y + block.end.y; y += blockDim.y) {
-    for (int x = threadIdx.x + block.start.x; x < threadIdx.x + block.end.x; x += blockDim.x) {
+  for (int y = threadIdx.y + block.start.y; y < block.end.y; y += blockDim.y) {
+    for (int x = threadIdx.x + block.start.x; x < block.end.x; x += blockDim.x) {
       out[y * sample.out_pitch.x + x] = ConvertSat<OutputType>(
               in[y * sample.in_pitch.x + x] * sample.contrast + sample.brightness);
     }

diff --git a/dali/operators/color/brightness_contrast.cc b/dali/operators/color/brightness_contrast.cc
@@ -66,6 +66,7 @@ bool BrightnessContrastCpu::SetupImpl(std::vector<OutputDesc> &output_desc,
 void BrightnessContrastCpu::RunImpl(workspace_t<CPUBackend> &ws) {
   const auto &input = ws.template InputRef<CPUBackend>(0);
   auto &output = ws.template OutputRef<CPUBackend>(0);
+  output.SetLayout(InputLayout(ws, 0));
   auto& tp = ws.GetThreadPool();
   TYPE_SWITCH(input.type().id(), type2id, InputType, (uint8_t, int16_t, int32_t, float), (
       TYPE_SWITCH(output_type_, type2id, OutputType, (uint8_t, int16_t, int32_t, float), (
@@ -86,5 +87,4 @@ void BrightnessContrastCpu::RunImpl(workspace_t<CPUBackend> &ws) {
   tp.WaitForWork();
 }
 
-
 }  // namespace dali
diff --git a/dali/operators/color/brightness_contrast.cu b/dali/operators/color/brightness_contrast.cu
@@ -52,6 +52,7 @@ bool BrightnessContrastGpu::SetupImpl(std::vector<OutputDesc> &output_desc,
 void BrightnessContrastGpu::RunImpl(workspace_t<GPUBackend> &ws) {
   const auto &input = ws.template Input<GPUBackend>(0);
   auto &output = ws.template Output<GPUBackend>(0);
+  output.SetLayout(InputLayout(ws, 0));
   TYPE_SWITCH(input.type().id(), type2id, InputType, (uint8_t, int16_t, int32_t, float), (
       TYPE_SWITCH(output_type_, type2id, OutputType, (uint8_t, int16_t, int32_t, float), (
           {

diff --git a/dali/operators/color/hsv.cc b/dali/operators/color/hsv.cc
@@ -24,7 +24,7 @@ using TheKernel = kernels::LinearTransformationCpu<Out, In, 3, 3, 3>;
 }  // namespace
 
 DALI_SCHEMA(Hsv)
-              .DocStr(R"code(This operator performs HSV manipulation.
+    .DocStr(R"code(This operator performs HSV manipulation.
 To change hue, saturation and/or value of the image, pass corresponding coefficients.
 Keep in mind, that `hue` has additive delta argument,
 while for `saturation` and `value` they are multiplicative.
@@ -36,18 +36,19 @@ The color vector is projected along the neutral (gray) axis,
 rotated (according to hue delta) and scaled according to value and saturation multiplers,
 and then restored to original color space.
 )code")
-              .NumInput(1)
-              .NumOutput(1)
-              .AddOptionalArg(hsv::kHue,
-                              R"code(Set additive change of hue. 0 denotes no-op)code", .0f,
-                              true)
-              .AddOptionalArg(hsv::kSaturation,
-                              R"code(Set multiplicative change of saturation. 1 denotes no-op)code",
-                              1.f, true)
-              .AddOptionalArg(hsv::kValue,
-                              R"code(Set multiplicative change of value. 1 denotes no-op)code",
-                              1.f, true)
-              .AddOptionalArg(hsv::kOutputType, R"code(Set output data type)code", DALI_UINT8);
+    .NumInput(1)
+    .NumOutput(1)
+    .AddOptionalArg(hsv::kHue,
+                    R"code(Set additive change of hue. 0 denotes no-op)code",
+                    0.0f, true)
+    .AddOptionalArg(hsv::kSaturation,
+                    R"code(Set multiplicative change of saturation. 1 denotes no-op)code",
+                    1.0f, true)
+    .AddOptionalArg(hsv::kValue,
+                    R"code(Set multiplicative change of value. 1 denotes no-op)code",
+                    1.0f, true)
+    .AddOptionalArg(hsv::kOutputType, R"code(Set output data type)code", DALI_UINT8)
+    .InputLayout(0, "HWC");
 
 
 DALI_REGISTER_OPERATOR(Hsv, HsvCpu, CPU)
@@ -77,6 +78,7 @@ bool HsvCpu::SetupImpl(std::vector<OutputDesc> &output_desc, const workspace_t<C
 void HsvCpu::RunImpl(workspace_t<CPUBackend> &ws) {
   const auto &input = ws.template InputRef<CPUBackend>(0);
   auto &output = ws.template OutputRef<CPUBackend>(0);
+  output.SetLayout(InputLayout(ws, 0));
   auto &tp = ws.GetThreadPool();
   TYPE_SWITCH(input.type().id(), type2id, InputType, (uint8_t, int16_t, int32_t, float, float16), (
       TYPE_SWITCH(output_type_, type2id, OutputType, (uint8_t, int16_t, int32_t, float, float16), (

diff --git a/dali/operators/color/hsv.cu b/dali/operators/color/hsv.cu
@@ -51,6 +51,7 @@ bool HsvGpu::SetupImpl(std::vector<OutputDesc> &output_desc, const workspace_t<G
 void HsvGpu::RunImpl(workspace_t<GPUBackend> &ws) {
   const auto &input = ws.template Input<GPUBackend>(0);
   auto &output = ws.template Output<GPUBackend>(0);
+  output.SetLayout(InputLayout(ws, 0));
   TYPE_SWITCH(input.type().id(), type2id, InputType, (uint8_t, int16_t, int32_t, float), (
       TYPE_SWITCH(output_type_, type2id, OutputType, (uint8_t, int16_t, int32_t, float), (
           {