aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorGeorgios Pinitas <georgios.pinitas@arm.com>2017-10-27 13:37:16 +0100
committerAnthony Barbier <anthony.barbier@arm.com>2018-11-02 16:35:24 +0000
commit1f378ee7639837601f5152ee9f185ee1a528d643 (patch)
treeba0b05fb5001c2e630c713f588b928bb51d97c3a
parent407c3e6e383affa7ebe72ce5f50fcf163ff037a3 (diff)
downloadComputeLibrary-1f378ee7639837601f5152ee9f185ee1a528d643.tar.gz
COMPMID-556: Static helper function for getting gws from window.
Change-Id: I3e47385b2f7a7977c2c61272f2fb30e800f39406 Reviewed-on: http://mpd-gerrit.cambridge.arm.com/93429 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>
-rw-r--r--arm_compute/core/CL/ICLKernel.h7
-rw-r--r--src/core/CL/ICLKernel.cpp23
2 files changed, 25 insertions, 5 deletions
diff --git a/arm_compute/core/CL/ICLKernel.h b/arm_compute/core/CL/ICLKernel.h
index 9119940bc5..a1bc3eb8d2 100644
--- a/arm_compute/core/CL/ICLKernel.h
+++ b/arm_compute/core/CL/ICLKernel.h
@@ -180,6 +180,13 @@ public:
* @return The maximum workgroup size value.
*/
size_t get_max_workgroup_size();
+ /** Get the global work size given an execution window
+ *
+ * @param[in] window Execution window
+ *
+ * @return Global work size of the given execution window
+ */
+ static cl::NDRange gws_from_window(const Window &window);
private:
/** Add the passed array's parameters to the object's kernel's arguments starting from the index idx.
diff --git a/src/core/CL/ICLKernel.cpp b/src/core/CL/ICLKernel.cpp
index 17b58b727f..13037a771d 100644
--- a/src/core/CL/ICLKernel.cpp
+++ b/src/core/CL/ICLKernel.cpp
@@ -43,15 +43,14 @@ void arm_compute::enqueue(cl::CommandQueue &queue, ICLKernel &kernel, const Wind
return;
}
- if((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
+ cl::NDRange gws = ICLKernel::gws_from_window(window);
+
+ // Check for empty NDRange
+ if(gws.dimensions() == 0)
{
return;
}
- cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
- (window.y().end() - window.y().start()) / window.y().step(),
- (window.z().end() - window.z().start()) / window.z().step());
-
cl::NDRange valid_lws;
if(lws_hint[0] * lws_hint[1] * lws_hint[2] > kernel.get_max_workgroup_size())
{
@@ -182,3 +181,17 @@ size_t ICLKernel::get_max_workgroup_size()
}
return _max_workgroup_size;
}
+
+cl::NDRange ICLKernel::gws_from_window(const Window &window)
+{
+ if((window.x().end() - window.x().start()) == 0 || (window.y().end() - window.y().start()) == 0)
+ {
+ return cl::NullRange;
+ }
+
+ cl::NDRange gws((window.x().end() - window.x().start()) / window.x().step(),
+ (window.y().end() - window.y().start()) / window.y().step(),
+ (window.z().end() - window.z().start()) / window.z().step());
+
+ return gws;
+} \ No newline at end of file