[SYCLomatic] modify image help function cases to support cuda backend.

Author: ziranzha

help_function/src/image_wrapper.cpp

@@ -12,19 +12,19 @@
 #include <dpct/dpct.hpp>
 
 dpct::image_wrapper<sycl::float4, 2> tex42;
-dpct::image_wrapper<sycl::char4, 1> tex_char4;
-dpct::image_wrapper<sycl::ushort4, 3> tex43;
+dpct::image_wrapper<sycl::int4, 1> tex_int4;
+dpct::image_wrapper<sycl::uint4, 3> tex43;
 
 // test_feature:image_accessor_ext
 void test_image(sycl::float4* out, dpct::image_accessor_ext<sycl::float4, 2> acc42,
-                  dpct::image_accessor_ext<sycl::char4, 1> acc21,
-                  dpct::image_accessor_ext<sycl::ushort4, 3> acc13) {
+                  dpct::image_accessor_ext<sycl::int4, 1> acc21,
+                  dpct::image_accessor_ext<sycl::uint4, 3> acc13) {
   // test_feature:read
   out[0] = acc42.read(0.5f, 0.5f);
   // test_feature:read
-  sycl::ushort4 data13 = acc13.read((uint64_t)3, (int)4, (short)5);
+  sycl::uint4 data13 = acc13.read((uint64_t)3, (int)4, (short)5);
   // test_feature:read
-  sycl::char4 data21 = acc21.read(0.5f);
+  sycl::int4 data21 = acc21.read(0.5f);
   int a = data21.x();
   int b = data13.x();
   out[1] = {a,b,a,b};
@@ -33,21 +33,21 @@ void test_image(sycl::float4* out, dpct::image_accessor_ext<sycl::float4, 2> acc
 int main() {
 
   sycl::float4 *host_buffer = new sycl::float4[640 * 480 * 24];
-  sycl::char4 *host_char_buffer = new sycl::char4[640];
+  sycl::int4 *host_char_buffer = new sycl::int4[640];
 
   for(int i = 0; i < 640 * 480 * 24; ++i) {
 	  host_buffer[i] = sycl::float4{10.0f, 10.0f, 10.0f, 10.0f};
 	  if (i < 640)
-		  host_char_buffer[i] = sycl::char4{30, 30, 30, 30};
+		  host_char_buffer[i] = sycl::int4{30, 30, 30, 30};
   }
   sycl::float4 *device_buffer;
   device_buffer = (sycl::float4 *)dpct::dpct_malloc(
                       640 * 480 * 24 * sizeof(sycl::float4));
   dpct::dpct_memcpy(device_buffer, host_buffer, 640 * 480 * 24 * sizeof(sycl::float4));
 
   dpct::image_channel chn1 =
-      dpct::image_channel(16, 16, 16, 16, dpct::image_channel_data_type::unsigned_int);
-  dpct::image_channel chn2 = dpct::image_channel::create<sycl::char4>();
+      dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::unsigned_int);
+  dpct::image_channel chn2 = dpct::image_channel::create<sycl::int4>();
   dpct::image_channel chn4 =
       dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::fp);
   chn4.set_channel_size(4, 32);
@@ -62,9 +62,9 @@ int main() {
 
   tex42.attach(image_data2, 640, 480, 650 * sizeof(sycl::float4));
 
-  dpct::dpct_memcpy(array1->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(host_buffer, 640 * sizeof(sycl::char4), 640 * sizeof(sycl::char4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * sizeof(sycl::char4), 1, 1));
+  dpct::dpct_memcpy(array1->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(host_buffer, 640 * sizeof(sycl::int4), 640 * sizeof(sycl::int4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * sizeof(sycl::int4), 1, 1));
   dpct::dpct_memcpy(dpct::pitched_data(image_data2, 650 * sizeof(sycl::float4), 640 * sizeof(sycl::float4*), 480), sycl::id<3>(0, 0, 0), dpct::pitched_data(host_buffer, 640 * 480 * sizeof(sycl::float4), 640 * 480 * sizeof(sycl::float4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * 480 * sizeof(sycl::float4), 1, 1));
-  dpct::dpct_memcpy(array3->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(device_buffer, 640 * 480 * 24 * sizeof(sycl::ushort4), 640 * 480 * 24 * sizeof(sycl::ushort4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * 480 * 24 * sizeof(sycl::ushort4), 1, 1));
+  dpct::dpct_memcpy(array3->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(device_buffer, 640 * 480 * 24 * sizeof(sycl::uint4), 640 * 480 * 24 * sizeof(sycl::uint4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * 480 * 24 * sizeof(sycl::uint4), 1, 1));
 
   tex43.attach(array3);
 
@@ -102,7 +102,7 @@ int main() {
     dpct::get_default_queue().submit([&](sycl::handler &cgh) {
       auto acc42 = tex42.get_access(cgh);
       auto acc13 = tex43.get_access(cgh);
-      auto acc21 = static_cast<dpct::image_wrapper<sycl::char4, 1> *>(tex41)->get_access(cgh);
+      auto acc21 = static_cast<dpct::image_wrapper<sycl::int4, 1> *>(tex41)->get_access(cgh);
 
       auto smpl42 = tex42.get_sampler();
       auto smpl13 = tex43.get_sampler();
@@ -112,8 +112,8 @@ int main() {
 
       cgh.single_task<dpct_kernel_name<class dpct_single_kernel>>([=] {
         test_image(acc_out.get_pointer(),dpct::image_accessor_ext<sycl::float4, 2>(smpl42, acc42),
-                   dpct::image_accessor_ext<sycl::char4, 1>(smpl21, acc21),
-                   dpct::image_accessor_ext<sycl::ushort4, 3>(smpl13, acc13));
+                   dpct::image_accessor_ext<sycl::int4, 1>(smpl21, acc21),
+                   dpct::image_accessor_ext<sycl::uint4, 3>(smpl13, acc13));
       });
     });
   }

help_function/src/image_wrapper_usm.cpp

@@ -12,13 +12,13 @@
 
 dpct::image_wrapper<sycl::float4, 2> tex42;
 dpct::image_wrapper<sycl::float4, 1> tex41;
-dpct::image_wrapper<sycl::ushort4, 3> tex43;
+dpct::image_wrapper<sycl::uint4, 3> tex43;
 
 void test_image(sycl::float4* out, dpct::image_accessor_ext<sycl::float4, 2> acc42,
                   dpct::image_accessor_ext<sycl::float4, 1> acc21,
-                  dpct::image_accessor_ext<sycl::ushort4, 3> acc13) {
+                  dpct::image_accessor_ext<sycl::uint4, 3> acc13) {
   out[0] = acc42.read(0.5f, 0.5f);
-  sycl::ushort4 data13 = acc13.read(0.5f, 0.5f, 0.5f);
+  sycl::uint4 data13 = acc13.read(0.5f, 0.5f, 0.5f);
   sycl::float4 data21 = acc21.read(0.5f);
   out[1].x() = data21.x();
   out[1].y() = data13.x();
@@ -37,7 +37,7 @@ int main() {
   dpct::dpct_memcpy(device_buffer, host_buffer, 640 * 480 * 24 * sizeof(sycl::float4));
 
   dpct::image_channel chn1 =
-      dpct::image_channel(16, 16, 16, 16, dpct::image_channel_data_type::unsigned_int);
+      dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::unsigned_int);
   dpct::image_channel chn2 =
       dpct::image_channel(32, 32, 32, 32, dpct::image_channel_data_type::fp);
   dpct::image_channel chn4 =
@@ -56,7 +56,7 @@ int main() {
 
   dpct::dpct_memcpy(array1->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(host_buffer, 640 * sizeof(sycl::float4), 640 * sizeof(sycl::float4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * sizeof(sycl::float4), 1, 1));
   dpct::dpct_memcpy(dpct::pitched_data(image_data2, 650 * sizeof(sycl::float4), 640 * sizeof(sycl::float4*), 480), sycl::id<3>(0, 0, 0), dpct::pitched_data(host_buffer, 640 * 480 * sizeof(sycl::float4), 640 * 480 * sizeof(sycl::float4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * 480 * sizeof(sycl::float4), 1, 1));
-  dpct::dpct_memcpy(array3->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(device_buffer, 640 * 480 * 24 * sizeof(sycl::ushort4), 640 * 480 * 24 * sizeof(sycl::ushort4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * 480 * 24 * sizeof(sycl::ushort4), 1, 1));
+  dpct::dpct_memcpy(array3->to_pitched_data(), sycl::id<3>(0, 0, 0), dpct::pitched_data(device_buffer, 640 * 480 * 24 * sizeof(sycl::uint4), 640 * 480 * 24 * sizeof(sycl::uint4), 1), sycl::id<3>(0, 0, 0), sycl::range<3>(640 * 480 * 24 * sizeof(sycl::uint4), 1, 1));
 
   tex43.attach(array3);
 
@@ -105,7 +105,7 @@ int main() {
       cgh.single_task<dpct_kernel_name<class dpct_single_kernel>>([=] {
         test_image(acc_out.get_pointer(),dpct::image_accessor_ext<sycl::float4, 2>(smpl42, acc42),
                    dpct::image_accessor_ext<sycl::float4, 1>(smpl21, acc21),
-                   dpct::image_accessor_ext<sycl::ushort4, 3>(smpl13, acc13));
+                   dpct::image_accessor_ext<sycl::uint4, 3>(smpl13, acc13));
       });
     });
   }