Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

OpenVX 1.3 - Laplacian Pyramid Node fix - GPU OpenCL #636

Merged
merged 36 commits into from
Oct 7, 2021
Merged

OpenVX 1.3 - Laplacian Pyramid Node fix - GPU OpenCL #636

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
36 commits
Select commit Hold shift + click to select a range
3edf373
replicate node fix
hansely123 Jun 11, 2021
4641fe5
replicate node fix
hansely123 Jun 11, 2021
978aaeb
replicate node
hansely123 Jun 15, 2021
e9f945e
Merge remote-tracking branch 'upstream/master' into H/replicatenode-gpu
hansely123 Jun 15, 2021
af5a9f5
fix
hansely123 Jun 15, 2021
11d3681
code cleanup
hansely123 Jun 15, 2021
fdb8446
code cleanup
hansely123 Jun 15, 2021
302cce2
Merge remote-tracking branch 'upstream/master' into H/replicatenode-gpu
hansely123 Jun 24, 2021
72b9f83
Merge branch 'GPUOpen-ProfessionalCompute-Libraries:master' into H/re…
hansely Jul 1, 2021
21434d7
warp affine bug fix (#9)
hansely Jul 6, 2021
b0db587
H/warpaffine gpu (#10)
hansely Jul 7, 2021
720f070
Merge branch 'GPUOpen-ProfessionalCompute-Libraries:master' into H/re…
hansely Jul 7, 2021
33dd0a8
warp affine fix for hip (#11)
hansely Jul 9, 2021
2d54789
Merge branch 'GPUOpen-ProfessionalCompute-Libraries:master' into H/re…
hansely Aug 25, 2021
aea6a4f
Laplacian
hansely123 Aug 25, 2021
2d4d943
Merge branch 'H/replicatenode-gpu' of https://github.com/hansely/MIVi…
hansely123 Aug 25, 2021
2604009
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Aug 30, 2021
f03057c
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Sep 3, 2021
40bddfe
temporary commit
hansely123 Sep 3, 2021
a85a80f
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Sep 10, 2021
e00a830
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Sep 13, 2021
20575fa
fix in progress
hansely123 Sep 15, 2021
59f39af
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Sep 15, 2021
edcb716
replace gaussian done
hansely123 Sep 16, 2021
4849244
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Sep 29, 2021
3e36158
pad first & last row of gpu kernel
hansely123 Sep 29, 2021
29e801f
code cleanup
hansely123 Sep 29, 2021
afc8e22
code cleanup
hansely123 Sep 29, 2021
85af7c3
code cleanup
hansely123 Sep 29, 2021
c8089ed
Merge remote-tracking branch 'upstream/master' into H/laplacian
hansely123 Sep 30, 2021
bd15cb4
pad first & last row for hip kernel
hansely123 Sep 30, 2021
09a3d84
kernel optimization
hansely123 Oct 7, 2021
66ea6ca
Merge branch 'GPUOpen-ProfessionalCompute-Libraries:master' into H/la…
hansely Oct 7, 2021
702e901
code cleanup
hansely123 Oct 7, 2021
07e8283
Merge branch 'H/laplacian' of https://github.com/hansely/MIVisionX in…
hansely123 Oct 7, 2021
acc796f
initialize to 0
hansely123 Oct 7, 2021
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
3 changes: 1 addition & 2 deletions amd_openvx/openvx/ago/ago_haf_cpu_pyramid.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -403,5 +403,4 @@ int HafCpu_ScaleGaussianOrb_U8_U8_5x5
pDstImage += dstImageStrideInBytes;
}
return AGO_SUCCESS;
}

}
5 changes: 3 additions & 2 deletions amd_openvx/openvx/ago/ago_haf_gpu_special_filters.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1403,7 +1403,7 @@ int HafGpu_ScaleGaussianHalf(AgoNode * node)
int work_group_width = 16;
int work_group_height = 16;
int width = node->paramList[0]->u.img.width;
int height = node->paramList[0]->u.img.height;
int height = node->paramList[0]->u.img.height - 1;
int N = 0;
if (node->akernel->id == VX_KERNEL_AMD_SCALE_GAUSSIAN_HALF_U8_U8_3x3) {
N = 3;
Expand Down Expand Up @@ -1437,6 +1437,7 @@ int HafGpu_ScaleGaussianHalf(AgoNode * node)
" int gx = get_global_id(0);\n"
" int gy = get_global_id(1);\n"
" p0_buf += p0_offset + (gy * p0_stride) + (gx << 2);\n"
" *(__global uint *)p0_buf = 0;\n"
" int gstride = p1_stride;\n"
" __global uchar * gbuf = p1_buf + p1_offset + (((gy - ly) << 1) + 1) * gstride + ((gx - lx) << 3);\n"
" bool valid = ((gx < %d) && (gy < %d)) ? true : false;\n" // (width+3)/4, height
Expand Down Expand Up @@ -1540,7 +1541,7 @@ int HafGpu_ScaleGaussianHalf(AgoNode * node)
" L0.s01 = vload2(0, (__local uint *)&lbuf_ptr[%d]);\n" // LMemStride * 4
" sum.s0 += (float)(L0.s0 & 0xffff); sum.s1 += (float)(L0.s0 >> 16); sum.s2 += (float)(L0.s1 & 0xffff); sum.s3 += (float)(L0.s1 >> 16);\n"
" sum = sum * (float4)0.00390625f;\n"
" if (valid) {;\n"
" if (valid) {\n"
" *(__global uint *)p0_buf = amd_pack(sum);\n"
" }\n"
"}\n"
Expand Down
3 changes: 2 additions & 1 deletion amd_openvx/openvx/ago/ago_kernel_api.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15493,8 +15493,9 @@ int agoKernel_ScaleGaussianHalf_U8_U8_5x5(AgoNode * node, AgoKernelCommand cmd)
status = VX_SUCCESS;
AgoData * oImg = node->paramList[0];
AgoData * iImg = node->paramList[1];
hipMemset(oImg->hip_memory, 0, oImg->size + oImg->gpu_buffer_offset);
kiritigowda marked this conversation as resolved.
Show resolved Hide resolved
if (HipExec_ScaleGaussianHalf_U8_U8_5x5(
node->hip_stream0, oImg->u.img.width, oImg->u.img.height,
node->hip_stream0, oImg->u.img.width, oImg->u.img.height - 1,
oImg->hip_memory + oImg->gpu_buffer_offset,oImg->u.img.stride_in_bytes,
iImg->u.img.width, iImg->u.img.height,
iImg->hip_memory + iImg->gpu_buffer_offset, iImg->u.img.stride_in_bytes, iImg->size)) {
Expand Down