Skip to content

Commit

Permalink
Update
Browse files Browse the repository at this point in the history
Signed-off-by: Jiang, Zhiwei <[email protected]>
  • Loading branch information
zhiweij1 committed May 14, 2024
1 parent e2447f2 commit a59d4a8
Show file tree
Hide file tree
Showing 6 changed files with 35 additions and 21 deletions.
8 changes: 3 additions & 5 deletions docs/dev_guide/reference/diagnostic_ref/dpct1122.rst
Original file line number Diff line number Diff line change
Expand Up @@ -37,11 +37,11 @@ results in the following migrated SYCL code:
.. code-block:: cpp
:linenos:
#include <sycl/sycl.hpp>
#include <dpct/dpct.hpp>
template <typename T> struct C { T a; };
/*
DPCT1122:0: 'longlong4' is migrated to 'sycl::long4' in the template declaration; it may cause template function or class redefinition; please adjust the code.
DPCT1122:0: 'longlong4' is migrated to 'sycl::long4' in the template
declaration; it may cause template function or class redefinition; please adjust
the code.
*/
template <> struct C<sycl::long4> { sycl::long4 a; };
template <> struct C<sycl::long4> { sycl::long4 a; };
Expand All @@ -51,8 +51,6 @@ which needs to be rewritten to:
.. code-block:: cpp
:linenos:
#include <sycl/sycl.hpp>
#include <dpct/dpct.hpp>
template <typename T> struct C { T a; };
template <> struct C<sycl::long4> { sycl::long4 a; };
6 changes: 4 additions & 2 deletions docs/dev_guide/reference/diagnostic_ref/dpct1123.rst
Original file line number Diff line number Diff line change
Expand Up @@ -62,7 +62,10 @@ results in the following migrated SYCL code:
kernel_array[10] = (void *)&kernel;
void *args[1] = {&d};
/*
DPCT1123:0: The kernel function pointer cannot be used in the device code. You need to call the kernel function with the correct argument(s) directly. According to the kernel function definition, adjusting the dimension of the sycl::nd_item may also be required.
DPCT1123:0: The kernel function pointer cannot be used in the device code. You
need to call the kernel function with the correct argument(s) directly.
According to the kernel function definition, adjusting the dimension of the
sycl::nd_item may also be required.
*/
q_ct1.parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, 16) * sycl::range<3>(1, 1, 16),
Expand All @@ -72,7 +75,6 @@ results in the following migrated SYCL code:
});
}
which needs to be rewritten to:

.. code-block:: cpp
Expand Down
13 changes: 6 additions & 7 deletions docs/dev_guide/reference/diagnostic_ref/dpct1124.rst
Original file line number Diff line number Diff line change
Expand Up @@ -38,18 +38,17 @@ results in the following migrated SYCL\* code:
:linenos:
/*
DPCT1124:0: cudaMemcpyAsync is migrated to asynchronous memcpy API. While the origin API might be synchronous, depends on the type of operand memory, so you may need to call wait() on event return by memcpy API to ensure synchronization behavior.
DPCT1124:0: cudaMemcpyAsync is migrated to asynchronous memcpy API. While the
origin API might be synchronous, it depends on the type of operand memory, so
you may need to call wait() on event return by memcpy API to ensure
synchronization behavior.
*/
dpct::device_ext &dev_ct1 = dpct::get_current_device();
sycl::queue &q_ct1 = dev_ct1.in_order_queue();
q_ct1.memcpy(host_dst, host_src, size);
dpct::get_in_order_queue().memcpy(host_dst, host_src, size);
which needs to be rewritten to:

.. code-block:: cpp
:linenos:
dpct::device_ext &dev_ct1 = dpct::get_current_device();
sycl::queue &q_ct1 = dev_ct1.in_order_queue();
q_ct1.memcpy(host_dst, host_src, size).wait();
dpct::get_in_order_queue().memcpy(host_dst, host_src, size).wait();
18 changes: 14 additions & 4 deletions docs/dev_guide/reference/diagnostic_ref/dpct1125.rst
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,10 @@ results in the following migrated SYCL code:
};
/*
DPCT1125:1: The type "Tk" defined in function "device" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the definition location of the type.
DPCT1125:1: The type "Tk" defined in function "device" is used as the parameter
type in all functions in the call path from the corresponding
sycl::handler::parallel_for() to the current function. You may need to adjust
the definition location of the type.
*/
template <typename T> void device(int a, int b, Tk *mem) {
using Tk = typename kernel_type_t<T>::Type;
Expand All @@ -67,7 +70,10 @@ results in the following migrated SYCL code:
}
/*
DPCT1125:2: The type "Tk" defined in function "device" is used as the parameter type in all functions in the call path from the corresponding sycl::handler::parallel_for() to the current function. You may need to adjust the definition location of the type.
DPCT1125:2: The type "Tk" defined in function "device" is used as the parameter
type in all functions in the call path from the corresponding
sycl::handler::parallel_for() to the current function. You may need to adjust
the definition location of the type.
*/
template <typename T>
void kernel(int a, int b, Tk *mem) {
Expand All @@ -77,7 +83,10 @@ results in the following migrated SYCL code:
template <typename T> void foo() {
int i;
/*
DPCT1126:0: The type "Tk" defined in function "device" is used as the parameter type in all functions in the call path from the sycl::handler::parallel_for() to the function "device". You may need to adjust the definition location of the type.
DPCT1126:0: The type "Tk" defined in function "device" is used as the
parameter type in all functions in the call path from the
sycl::handler::parallel_for() to the function "device". You may need to adjust
the definition location of the type.
*/
dpct::get_in_order_queue().submit([&](sycl::handler &cgh) {
sycl::local_accessor<Tk, 1> mem_acc_ct1(sycl::range<1>(256), cgh);
Expand All @@ -87,7 +96,8 @@ results in the following migrated SYCL code:
[=](sycl::nd_item<3> item_ct1) {
kernel<T>(
i, i,
mem_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>().get());
mem_acc_ct1.template get_multi_ptr<sycl::access::decorated::no>()
.get());
});
});
}
Expand Down
4 changes: 3 additions & 1 deletion docs/dev_guide/reference/diagnostic_ref/dpct1127.rst
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,9 @@ results in the following migrated SYCL code:
// test.dp.cpp
/*
DPCT1127:0: The constant compile-time initialization for device_global is supported when compiling with C++20. You may need to adjust the compile commands.
DPCT1127:0: The constant compile-time initialization for device_global is
supported when compiling with C++20. You may need to adjust the compile
commands.
*/
static sycl::ext::oneapi::experimental::device_global<const int> var{0};
void kernel(int* ptr) {
Expand Down
7 changes: 5 additions & 2 deletions docs/dev_guide/reference/diagnostic_ref/dpct1128.rst
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,8 @@ results in the following migrated SYCL code:
/*
DPCT1128:0: The type "UserStruct<float, int>" is not device copyable for copy
constructor breaking the device copyable requirement. It is used in the SYCL kernel, please rewrite the code.
constructor breaking the device copyable requirement. It is used in the SYCL
kernel, please rewrite the code.
*/
template <class T1, class T2> struct UserStruct {
UserStruct() {}
Expand All @@ -63,7 +64,9 @@ results in the following migrated SYCL code:
void foo() {
UserStruct<float, int> us;
/*
DPCT1129:1: The type "UserStruct<float, int>" is used in the SYCL kernel but it is not device copyable. The sycl::is_device_copyable specialization has been added for this type, please review the code.
DPCT1129:1: The type "UserStruct<float, int>" is used in the SYCL kernel, but
it is not device copyable. The sycl::is_device_copyable specialization has
been added for this type. Please review the code.
*/
dpct::get_in_order_queue().parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
Expand Down

0 comments on commit a59d4a8

Please sign in to comment.