Skip to content

Commit

Permalink
Release 6.3.1
Browse files Browse the repository at this point in the history
* Fixed Vulkan.Screenshot, so it captures it at the correct time.
* OpenCL.SoftISP example updated.
* When UI.Benchmark is using DrawCaching any configuration change will force invalidate the cache so the effect of the options is immediately visible.
* Windows did not get marked as dirty by layout changes.
* Updated zlib version
* OpenCL SoftISP Fix denoise kernel issue

Change-Id: I00f9a8d624905e6f2af43ca686d4ac763bc85ddd
  • Loading branch information
ReneThrane committed May 1, 2024
1 parent 91a9eb1 commit 4bdb1c4
Show file tree
Hide file tree
Showing 32 changed files with 1,943 additions and 1,719 deletions.
1 change: 0 additions & 1 deletion .clang-format
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,6 @@ ColumnLimit: '150'
ConstructorInitializerAllOnOneLineOrOnePerLine: false
ConstructorInitializerIndentWidth: '2'
ContinuationIndentWidth: '2'
Cpp11BracedListStyle: true
DerivePointerAlignment: false
EmptyLineBeforeAccessModifier: Always
EmptyLineAfterAccessModifier: Never
Expand Down
2 changes: 1 addition & 1 deletion CONTRIBUTING.md
Original file line number Diff line number Diff line change
Expand Up @@ -221,6 +221,6 @@ These are just recommended third party tools and plugins, but use them at your o
* [PlantUML](https://marketplace.visualstudio.com/items?itemName=jebbs.plantuml)
* [Python](https://marketplace.visualstudio.com/items?itemName=ms-python.python)
* [Visual Studio Keymap](https://marketplace.visualstudio.com/items?itemName=ms-vscode.vs-keybindings)
* [Visual Studio 2022](https://visualstudio.microsoft.com/)
* [Visual Studio 2022](https://visualstudio.microsoft.com/)
* [Markdown Editor](https://marketplace.visualstudio.com/items?itemName=MadsKristensen.MarkdownEditor)
* [Trailing Whitespace Visualizer](https://marketplace.visualstudio.com/items?itemName=MadsKristensen.TrailingWhitespaceVisualizer)
225 changes: 111 additions & 114 deletions DemoApps/GLES3/OpenCLGaussianFilter/Content/gaussian_filter.cl
Original file line number Diff line number Diff line change
Expand Up @@ -12,119 +12,116 @@
* Gaussian OpenCL (Vivante) Kernel
*/

__kernel void gaussian_filter ( __global uchar *input,
__global uchar *output,
int width,
int height,
int channels

)
__kernel void gaussian_filter(__global uchar* input, __global uchar* output, int width, int height, int channels

)
{
int x = get_global_id (0);
int y = get_global_id (1);
int id = (y * width) + x;
int id1 = (y * width) + x+1;
int id2 = (y * width) + x+2;

float sum = 0;
float sum2 = 0;
float sum3 = 0;

float kernel_weights[3][3] = {
{1.0/16.0, 2.0/16.0, 1.0/16.0},
{2.0/16.0, 4.0/16.0, 2.0/16.0},
{1.0/16.0, 2.0/16.0, 1.0/16.0}
};

if ((id - width -1) < 0 || (id + width + 1) >= (width * height - 1))
{
output[id] = input[id];
return;
}

if (channels == 3)
{
if ((id1 - width -1) < 0 || (id1 + width + 1) >= (width * height - 1))
{
output[id] = input[id];
return;
}
if ((id2 - width -1) < 0 || (id2 + width + 1) >= (width * height - 1))
{
output[id] = input[id];
return;
}
}

if (channels == 1)
{
sum = input[id - width - 1] * kernel_weights[0][0];
sum += input[id - width] * kernel_weights[0][1];
sum += input[id - width + 1] * kernel_weights[0][2];

sum += input[id - 1] * kernel_weights[1][0];
sum += input[id] * kernel_weights[1][1];
sum += input[id + 1] * kernel_weights[1][2];

sum += input[id + width - 1] * kernel_weights[2][0];
sum += input[id + width] * kernel_weights[2][1];
sum += input[id + width + 1] * kernel_weights[2][2];

output[id] = sum;
}

else if (channels == 3)
{

// R Channel
sum = input[id - width - 1] * kernel_weights[0][0];
sum += input[id - width] * kernel_weights[0][1];
sum += input[id - width + 1] * kernel_weights[0][2];

sum += input[id - 1] * kernel_weights[1][0];
sum += input[id] * kernel_weights[1][1];
sum += input[id + 1] * kernel_weights[1][2];

sum += input[id + width - 1] * kernel_weights[2][0];
sum += input[id + width] * kernel_weights[2][1];
sum += input[id + width + 1] * kernel_weights[2][2];

// G Channel
sum2 = input[id1 - width - 1] * kernel_weights[0][0];
sum2 += input[id1 - width] * kernel_weights[0][1];
sum2 += input[id1 - width + 1] * kernel_weights[0][2];

sum2 += input[id1 - 1] * kernel_weights[1][0];
sum2 += input[id1] * kernel_weights[1][1];
sum2 += input[id1 + 1] * kernel_weights[1][2];

sum2 += input[id1 + width - 1] * kernel_weights[2][0];
sum2 += input[id1 + width] * kernel_weights[2][1];
sum2 += input[id1 + width + 1] * kernel_weights[2][2];


// B Channel
sum3 = input[id2 - width - 1] * kernel_weights[0][0];
sum3 += input[id2 - width] * kernel_weights[0][1];
sum3 += input[id2 - width + 1] * kernel_weights[0][2];

sum3 += input[id2 - 1] * kernel_weights[1][0];
sum3 += input[id2] * kernel_weights[1][1];
sum3 += input[id2 + 1] * kernel_weights[1][2];

sum3 += input[id2 + width - 1] * kernel_weights[2][0];
sum3 += input[id2 + width] * kernel_weights[2][1];
sum3 += input[id2 + width + 1] * kernel_weights[2][2];

output[id] = sum;
output[id1] = sum2;
output[id2] = sum3;
}

else
{
output[id] = 0;
output[id1] = 255;
output[id2] = 255;
}
int x = get_global_id(0);
int y = get_global_id(1);
int id = (y * width) + x;
int id1 = (y * width) + x + 1;
int id2 = (y * width) + x + 2;

float sum = 0;
float sum2 = 0;
float sum3 = 0;

float kernel_weights[3][3] = {
// clang-format off
{1.0/16.0, 2.0/16.0, 1.0/16.0},
{2.0/16.0, 4.0/16.0, 2.0/16.0},
{1.0/16.0, 2.0/16.0, 1.0/16.0}
// clang-format on
};

if ((id - width - 1) < 0 || (id + width + 1) >= (width * height - 1))
{
output[id] = input[id];
return;
}

if (channels == 3)
{
if ((id1 - width - 1) < 0 || (id1 + width + 1) >= (width * height - 1))
{
output[id] = input[id];
return;
}
if ((id2 - width - 1) < 0 || (id2 + width + 1) >= (width * height - 1))
{
output[id] = input[id];
return;
}
}

if (channels == 1)
{
sum = input[id - width - 1] * kernel_weights[0][0];
sum += input[id - width] * kernel_weights[0][1];
sum += input[id - width + 1] * kernel_weights[0][2];

sum += input[id - 1] * kernel_weights[1][0];
sum += input[id] * kernel_weights[1][1];
sum += input[id + 1] * kernel_weights[1][2];

sum += input[id + width - 1] * kernel_weights[2][0];
sum += input[id + width] * kernel_weights[2][1];
sum += input[id + width + 1] * kernel_weights[2][2];

output[id] = sum;
}

else if (channels == 3)
{
// R Channel
sum = input[id - width - 1] * kernel_weights[0][0];
sum += input[id - width] * kernel_weights[0][1];
sum += input[id - width + 1] * kernel_weights[0][2];

sum += input[id - 1] * kernel_weights[1][0];
sum += input[id] * kernel_weights[1][1];
sum += input[id + 1] * kernel_weights[1][2];

sum += input[id + width - 1] * kernel_weights[2][0];
sum += input[id + width] * kernel_weights[2][1];
sum += input[id + width + 1] * kernel_weights[2][2];

// G Channel
sum2 = input[id1 - width - 1] * kernel_weights[0][0];
sum2 += input[id1 - width] * kernel_weights[0][1];
sum2 += input[id1 - width + 1] * kernel_weights[0][2];

sum2 += input[id1 - 1] * kernel_weights[1][0];
sum2 += input[id1] * kernel_weights[1][1];
sum2 += input[id1 + 1] * kernel_weights[1][2];

sum2 += input[id1 + width - 1] * kernel_weights[2][0];
sum2 += input[id1 + width] * kernel_weights[2][1];
sum2 += input[id1 + width + 1] * kernel_weights[2][2];


// B Channel
sum3 = input[id2 - width - 1] * kernel_weights[0][0];
sum3 += input[id2 - width] * kernel_weights[0][1];
sum3 += input[id2 - width + 1] * kernel_weights[0][2];

sum3 += input[id2 - 1] * kernel_weights[1][0];
sum3 += input[id2] * kernel_weights[1][1];
sum3 += input[id2 + 1] * kernel_weights[1][2];

sum3 += input[id2 + width - 1] * kernel_weights[2][0];
sum3 += input[id2 + width] * kernel_weights[2][1];
sum3 += input[id2 + width + 1] * kernel_weights[2][2];

output[id] = sum;
output[id1] = sum2;
output[id2] = sum3;
}

else
{
output[id] = 0;
output[id1] = 255;
output[id2] = 255;
}
}
Loading

0 comments on commit 4bdb1c4

Please sign in to comment.