remove unused 'out_stream' from SYCL variants (which improves the per…

…formance as well)
ecmwf-ifs · Jan 17, 2024 · e615fdf · e615fdf
1 parent 50bcf3f
commit e615fdf
Show file tree

Hide file tree

Showing 6 changed files with 6 additions and 9 deletions.
diff --git a/src/cloudsc_sycl/cloudsc/cloudsc_c.kernel b/src/cloudsc_sycl/cloudsc/cloudsc_c.kernel
@@ -39,7 +39,7 @@ void cloudsc_c(int kidia, int kfdia, int klon, double ptsphy,
   double r5ies, double r5alvcp, double r5alscp, double ralvdcp, double ralsdcp, double ralfdcp,
   double rtwat, double rtice, double rticecu, double rtwat_rtice_r, double rtwat_rticecu_r,
   double rkoop1, double rkoop2,
-  cl::sycl:: stream out_stream, cl::sycl::nd_item<1> item_ct) {
+  cl::sycl::nd_item<1> item_ct) {
 
   //-------------------------------------------------------------------------------
   //                 Declare input/output arguments

diff --git a/src/cloudsc_sycl/cloudsc/cloudsc_c_hoist.kernel b/src/cloudsc_sycl/cloudsc/cloudsc_c_hoist.kernel
@@ -45,7 +45,7 @@ void cloudsc_c(int kidia, int kfdia, int klon, double ptsphy,
   double * __restrict__ zpfplsx, double * __restrict__ zlneg, double * __restrict__ zqxn2d,
   double * __restrict__ zqsmix, double * __restrict__ zqsliq, double * __restrict__ zqsice,
   double * __restrict__ zfoeewmt, double * __restrict__ zfoeew, double * __restrict__ zfoeeliqt,
-  cl::sycl:: stream out_stream, cl::sycl::nd_item<1> item_ct) {
+  cl::sycl::nd_item<1> item_ct) {
 
   //-------------------------------------------------------------------------------
   //                 Declare input/output arguments

diff --git a/src/cloudsc_sycl/cloudsc/cloudsc_c_k_caching.kernel b/src/cloudsc_sycl/cloudsc/cloudsc_c_k_caching.kernel
@@ -39,7 +39,7 @@ void cloudsc_c(int kidia, int kfdia, int klon, double ptsphy,
   double r5ies, double r5alvcp, double r5alscp, double ralvdcp, double ralsdcp, double ralfdcp,
   double rtwat, double rtice, double rticecu, double rtwat_rtice_r, double rtwat_rticecu_r,
   double rkoop1, double rkoop2,
-  cl::sycl:: stream out_stream, cl::sycl::nd_item<1> item_ct) {
+  cl::sycl::nd_item<1> item_ct) {
 
   //-------------------------------------------------------------------------------
   //                 Declare input/output arguments

diff --git a/src/cloudsc_sycl/cloudsc/cloudsc_driver.cpp b/src/cloudsc_sycl/cloudsc/cloudsc_driver.cpp
@@ -398,7 +398,6 @@ void cloudsc_driver(int numthreads, int numcols, int nproma) {
     cl::sycl::range<1> local(nproma);
 
     q.submit([&](cl::sycl::handler &h) {
-        cl::sycl::stream out_stream(16384, 512, h);
         h.parallel_for( cl::sycl::nd_range<1>( global, local), [=] (cl::sycl::nd_item<1> item_ct1) {			 
 
     cloudsc_c(1, icend, nproma, ptsphy, d_pt, d_pq,
@@ -420,7 +419,7 @@ void cloudsc_driver(int numthreads, int numcols, int nproma) {
                 rv, r2es, r3les, r3ies, r4les, r4ies, r5les,
                 r5ies, r5alvcp, r5alscp, ralvdcp, ralsdcp, ralfdcp,
                 rtwat, rtice, rticecu, rtwat_rtice_r, rtwat_rticecu_r,
-                rkoop1, rkoop2, out_stream, item_ct1);
+                rkoop1, rkoop2, item_ct1);
 
 	});
     });

diff --git a/src/cloudsc_sycl/cloudsc/cloudsc_driver_hoist.cpp b/src/cloudsc_sycl/cloudsc/cloudsc_driver_hoist.cpp
@@ -431,7 +431,6 @@ void cloudsc_driver(int numthreads, int numcols, int nproma) {
     cl::sycl::range<1> local(nproma);
 
     q.submit([&](cl::sycl::handler &h) {
-        cl::sycl::stream out_stream(16384, 512, h);
         h.parallel_for( cl::sycl::nd_range<1>( global, local), [=] (cl::sycl::nd_item<1> item_ct1) {
 
     cloudsc_c(1, icend, nproma, ptsphy, d_pt, d_pq,
@@ -460,7 +459,7 @@ void cloudsc_driver(int numthreads, int numcols, int nproma) {
                 d_zpfplsx, d_zlneg, d_zqxn2d,
                 d_zqsmix, d_zqsliq, d_zqsice,
                 d_zfoeewmt, d_zfoeew, d_zfoeeliqt, 
-		out_stream, item_ct1);
+		item_ct1);
 
 
     });

diff --git a/src/cloudsc_sycl/cloudsc/cloudsc_driver_k_caching.cpp b/src/cloudsc_sycl/cloudsc/cloudsc_driver_k_caching.cpp
@@ -396,7 +396,6 @@ void cloudsc_driver(int numthreads, int numcols, int nproma) {
     cl::sycl::range<1> local(nproma);
 
     q.submit([&](cl::sycl::handler &h) {
-        cl::sycl::stream out_stream(16384, 512, h);
         h.parallel_for( cl::sycl::nd_range<1>( global, local), [=] (cl::sycl::nd_item<1> item_ct1) {			 
 
     cloudsc_c(1, icend, nproma, ptsphy, d_pt, d_pq,
@@ -418,7 +417,7 @@ void cloudsc_driver(int numthreads, int numcols, int nproma) {
                 rv, r2es, r3les, r3ies, r4les, r4ies, r5les,
                 r5ies, r5alvcp, r5alscp, ralvdcp, ralsdcp, ralfdcp,
                 rtwat, rtice, rticecu, rtwat_rtice_r, rtwat_rticecu_r,
-                rkoop1, rkoop2, out_stream, item_ct1);
+                rkoop1, rkoop2, item_ct1);
 
 	});
     });