Developer Guide and Reference

  • 2021.1
  • 12/04/2020
  • Public Content

Diagnostics Reference

This topic gives you all of the diagnostic reference numbers, messages, detailed help, and suggestion to fix your errors.

Diagnostics Reference

ID
Message
Detailed Help
Suggestions to Fix
DPCT1000
An error handling
if-stmt
was detected but could not be rewritten. See the details in the resulting file comments.
The CUDA* API return error codes that are consumed by the program logic. SYCL* uses exceptions to report errors and does not return the error code.
When the error handling logic in the original code is simple (for example, a print error message and exit), the code is removed in the resulting
Data Parallel C++ (
DPC++
)
application. The expectation is that SYCL throws an exception, which is handled with the printing of an exception message and exiting (the exception handler is generated automatically by the
Intel®
DPC++
Compatibility Tool
).
This warning is generated when the
Intel®
DPC++
Compatibility Tool
detects more complex error handling than it considers safe to remove.
Review the error handling
if-
statement and try to rewrite it to use an exception handler instead.
DPCT1001
The statement could not be removed. See the details in the resulting file comments.
The
Intel®
DPC++
Compatibility Tool
was not able to remove the code in the
then
clause of
if-stmt
. See
DPCT1000
.
See
DPCT1000
.
DPCT1002
A special case error handling
if-stmt
was detected. You may need to rewrite this code.
See
DPCT1000
See
DPCT1000
.
DPCT1003
The migrated API does not return an error code, so (*, 0) is inserted. You may need to rewrite this code.
Typically, this happens because the CUDA API returns an error code and then it is consumed by the program logic.
SYCL uses exceptions to report errors and does not return the error code.
The
Intel®
DPC++
Compatibility Tool
inserts a (*, 0) operator, so that the resulting application can be compiled. This operator returns 0 and is inserted if the return code is expected by the program logic and the new API does not return it. You should review all such places in the code.
If in a
DPC++
application you:
  • Do not need the code that consumes the error code, remove the code and the (*, 0) operator.
  • Need the code that consumes the error code, try to replace it with an exception handling code and use your logic in an exception handler.
DPCT1004
Could not generate a replacement.
Contact Support for help.
Report the issue.
DPCT1005
The device version is different. You may need to rewrite this code.
The
Intel®
DPC++
Compatibility Tool
detected a usage of CUDA Compute Capability-dependent logic in the original program.
DPC++
does not support CUDA Compute Capability. The logic in the generated code uses a version extracted from the
cl::sycl::info::device::version
, which is different than CUDA Compute Capability.
Review the logic and adjust it.
DPCT1006
DPC++
does not provide a standard API to differentiate between integrated/discrete GPU devices.
The generated code uses a
0
(zero) as the value of this field.
Consider re-implementing the code that depends on this field.
DPCT1007
Migration of this API is not supported by the
Intel®
DPC++
Compatibility Tool
.
Contact Support for help.
This kind of code requires manual migration to
DPC++
.
DPCT1008
The clock function is not defined in
DPC++
. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement.
The clock call was not replaced, because it is not defined by
DPC++
. Replace this hardware-specific feature with the one offered by your hardware vendor.
Consult with your hardware vendor to find a replacement.
DPCT1009
SYCL uses exceptions to report errors and does not use error codes. The original code was commented out and a warning string was inserted. You may need to rewrite this code.
SYCL uses exceptions to report errors and does not use error codes. The original code tries to get a string message through the error code, while SYCL does not require such functionality.
To indicate that the code needs to be updated, a warning string is inserted.
You may need to rewrite this code.
DPCT1010
SYCL uses exceptions to report errors and does not use error codes. The call was replaced with 0. You may need to rewrite this code.
SYCL uses exceptions to report errors and does not use error codes. The original code tries to get an error code, while SYCL does not require such functionality. The call was replaced with 0.
You may need to rewrite this code.
DPCT1011
The tool detected overloaded operators for built-in vector types, which may conflict with the SYCL 1.2.1 standard operators (see 4.10.2.1 Vec interface). The tool inserted a namespace to avoid the conflict. Use SYCL 1.2.1 standard operators instead.
For vector types, like
double2
, you may have overloaded operators, while overloaded operators with the same signature are also defined in SYCL standard, resulting in a conflict. The
Intel®
DPC++
Compatibility Tool
adds the namespace for overloaded operators overloaded to differentiate them from the ones defined in SYCL. You may need to rewrite the code.
You may need to rewrite this code.
DPCT1012
Detected a kernel execution time measurement pattern and generated an initial code for time measurements in SYCL. You can change the way time is measured depending on your requirements.
The generated code uses the CPU time to measure the kernel execution time. You can change the way time is measured depending on your requirements.
Review the logic and adjust it as needed.
DPCT1013
The rounding mode could not be specified and the generated code may have different precision than the original code. Verify the correctness. The SYCL math built-in functions rounding mode is aligned with the OpenCL™ C 1.2 standard.
The rounding mode could not be specified and the generated code may have different precision than the original code. Verify the correctness. The SYCL math built-in functions rounding mode is aligned with the OpenCL C 1.2 standard.
Review the logic and adjust it.
DPCT1014
The flag/priority options are not supported for SYCL queues; the output parameter(s) are set to 0.
The flag and priority options are not supported in the SYCL queue. You may need to rewrite the generated code.
Review the logic and adjust it.
DPCT1015
The output requires adjustment.
Use the SYCL output stream instead of a
printf
-like formatted string.
Use a SYCL output stream instead of a
printf
-like formatted string.
DPCT1016
The <API name> was not migrated, because parameter(s) <parameter name a> and/or <parameter name b> could not be evaluated, or because <parameter name a> is not equal to <parameter name b>. Rewrite this code manually.
The
cublasSetMatrix
can be migrated by the
Intel®
DPC++
Compatibility Tool
only when
lda
and
ldb
have the same constant value. In cases where the migration did not occur, rewrite the code manually:
  • If the values of
    lda
    and
    ldb
    are the same, you can use the following code:
    dpct::dpct_memcpy((void*)(B), (void*)(A), lda*cols*elemSize, dpct::host_to_device); // For cublasGetMatrix, use dpct::device_to_host
  • Otherwise, you can copy the columns of the matrix or elements of the vector one by one:
    • The equivalent of
      cublasSetMatrix(rows, cols, elemSize, A, lda, B, ldb)
      is:
      auto A_backup = A; auto B_backup = B; A = A - lda; B = B - ldb; for(int c = 0; c < cols; ++c) { A = A + lda; B = B + ldb; dpct::dpct_memcpy((void*)(B), (void*)(A), rows*elemSize, dpct::host_to_device); // For cublasGetMatrix, use dpct::device_to_host } A = A_backup; B = B_backup;
The
Intel®
DPC++
Compatibility Tool
can migrate the
cublasSetVector
only when
incx
and
incy
have the same constant value. In cases where the migration does not occur, rewrite the code manually:
  • If the values of
    incx
    and
    incy
    are the same, you can use the following code:
    dpct::dpct_memcpy((void*)(B), (void*)(A), n*incx*elemSize, dpct::host_to_device); // For cublasGetVector, use dpct::device_to_host
  • Otherwise, you can copy the elements of the vector one by one:
    • To replace
      cublasGetVector(n, elemSize, x, incx, y, incy)
      , use the following snippet:
      auto x_backup = x; auto y_backup = y; x = x - incx; y = y - incy; for(int c = 0; c < n; ++c) { x = x + incx; y = y + incy; dpct::dpct_memcpy((void*)(y), (void*)(x), elemSize, dpct::device_to_host); // for cublasSetVector, use dpct::host_to_device } x = x_backup; y = y_backup;
Review the logic and adjust it.
DPCT1017
The <DPC++ API name> call is used instead of the <CUDA API name> call. These two calls do not provide the same functionality. Check the potential precision and/or performance issues for the generated code.
The
Intel®
DPC++
Compatibility Tool
did not find an exact functional equivalent of the function in
DPC++
, so it used the closest alternative. The replacement may impact your code precision and/or performance. Verify the code correctness/performance.
Verify the code correctness/performance.
DPCT1018
The <API name> was migrated, but the generated code performance may be sub-optimal due to the following reasons:
  • The <parameter name> could not be evaluated and may be smaller than <other parameter name>.
  • The <parameter name> is equal to <other parameter name>, but is greater than 1.
This warning appears in two cases:
  1. cublasSetMatrix
    The
    Intel®
    DPC++
    Compatibility Tool
    replaced the
    cublasSetMatrix
    with memory copying from the host to the device. When the rows parameter of the
    cublasSetMatrix
    is smaller than the
    lda
    parameter, the generated code copies more data (
    lda*cols
    ) than the actual data available in the matrix (
    rows*cols
    ).
    To improve performance, consider changing the values of
    lda
    and
    ldb
    . If the rows parameter is greater than or equal to
    lda
    , no action is required for this code.
  2. cublasSetVector
    The
    Intel®
    DPC++
    Compatibility Tool
    replaced the
    cublasSetVector
    with memory copying from the host to the device. When the
    incx
    parameter of the
    cublasSetVector
    equals the
    incy
    parameter, but is greater than 1, the generated code copies more data (
    incx
    *
    n
    ) than the actual data available in the vector (
    n
    ). To improve performance, consider changing the values of
    incx
    and
    incy
    .
If the rows parameter of the
cublasSetMatrix
is smaller than the
lda
parameter and you observe performance issues, consider changing the values of
lda
and
ldb
.
If the
incx
parameter of the
cublasSetVector
equals the
incy
parameter, but is greater than 1 and you observe performance issues, consider changing the values of
incx
and
incy
.
DPCT1019
The
local_mem_size
in SYCL is not a complete equivalent of
sharedMemPerBlock
. You may need to adjust the code.
In CUDA, the
sharedMemPerBlock
reports the size of the shared memory in bytes available per block. The SYCL equivalent of a CUDA block is a work-group. The SYCL equivalent of shared memory is local memory. There is no limitation on the size of the local memory per work-group in SYCL. There is a limit on the maximum size of the local memory in bytes available per compute unit, which is exposed by the
info::device::local_mem_size
device descriptor in SYCL.
Verify the code correctness.
DPCT1020
The migration of <api name>, if it is called from
__global__
or
__device__
function is not supported. You may need to redesign the code to use the host-side <api name> instead, which submits this call to the SYCL queue automatically.
The warning message is generated in cases where the
<api name>
itself submits the SYCL kernel to the command queue, and the caller of
<api-name>
is the SYCL kernel that is submitted to the command queue itself. It results in device-side enqueue of the kernel, which is not supported by SYCL 1.2.1.
Redesign the code to use the host-side API, which submits this call to the SYCL queue automatically.
DPCT1021
Migration of
cublasHandle_t
in
__global__
or
__device__
function is not supported. You may need to redesign the code
The
cublasHandle_t
is migrated to the
cl::sycl::queue
. Because SYCL 1.2.1 does not support device-side enqueue of the kernel, the usage of the
cl::sycl::queue
is not possible in SYCL kernels.
Redesign the code to submit all kernels by the host program instead of device-side enqueue.
DPCT1022
There is no exact match between the
maxGridSize
and the
max_nd_range
size. Verify the correctness of the code.
There is no analogue of the
maxGridSize
in SYCL. SYCL
nd_ranges
can have up to three dimensions, just like grids in CUDA, but there is no maximum of
nd_range
size beyond the data type width, which is
size_t
. The
Intel®
DPC++
Compatibility Tool
replaces the
maxGridSize
with the
max_nd_range_size
helper, which is initialized to the
size_t
width.
Verify the code correctness.
DPCT1023
The
DPC++
sub-group does not support mask options for <api name>.
This warning is generated when the
DPC++
sub-group function used for migration does not have mask parameter.
Verify the code correctness.
DPCT1024
The original code returned the error code that was further consumed by the program logic. This original code was replaced with 0. You may need to rewrite the program logic consuming the error code.
This warning is generated in cases where in the original code, the CUDA API call returns the error code, which is consumed by the program logic:
handleError(cudaEventRecord(e));
If in the resulting code the CUDA API call is replaced by the code, which does not return the error code, 0 is used as an input to the program logic, consuming the error code:
e_ct1 = clock(), handleError(0);
The error handling code in that case must be verified and may require replacing it with exception handling code or removed completely.
Verify the code correctness.
DPCT1025
The SYCL queue is created ignoring the flag/priority options.
The flag and priority options are not supported when creating the SYCL queue. You may need to rewrite the generated code.
Verify the code correctness.
DPCT1026
The call to <API name> was removed, because <reason>.
The reason can be:
  • DPC++
    currently does not support setting resource limits on devices
  • DPC++
    currently does not support associating USM with a specific queue
  • DPC++
    currently does not support query operations on queues
  • DPC++
    currently does not support capture operations on queues
  • DPC++
    currently does not support configuring shared memory on devices
API calls from the original application, which do not have functionally compatible
DPC++
API calls are removed if the
Intel®
DPC++
Compatibility Tool
determines that it should not affect the program logic.
Verify the code correctness.
DPCT1027
The call to <API name> was replaced with 0, because <reason>.
The reason can be:
  • DPC++
    currently does not support setting resource limits on devices
  • DPC++
    currently does not support associating USM with a specific queue
  • DPC++
    currently does not support query operations on queues
  • DPC++
    currently does not support capture operations on queues
  • DPC++
    currently does not support configuring shared memory on devices
API calls from the original application, which do not have functionally compatible
DPC++
API calls are replaced with 0 if the
Intel®
DPC++
Compatibility Tool
determines that this call removal should not affect the program logic and at the same time the return value of that call is used for error handling.
Verify the code correctness.
DPCT1028
The
<API name>
was not migrated, because the parameter
<parameter name>
is unsupported.
The
DPC++
Compatibility Tool was not able to recognize the parameter value used as an engine type.
Rewrite this code manually by using a supported random number generator. See the Random Number Generators topic for more information.
DPCT1029
DPC++
currently does not support getting limits on devices. The output parameter(s) are set to 0.
DPC++
currently does not support getting limits on devices. The output parameter(s) are set to 0.
Contact Support for help.
Review the logic and adjust it manually.
DPCT1030
DPC++
currently does not support inter-process communication (IPC) operations. You may need to rewrite the code.
IPC operations are currently not supported by
DPC++
.
Contact Support for help.
Rewrite the code manually.
DPCT1031
DPC++
currently does not support memory access across peer devices. The output parameter(s) are set 0.
DPC++
currently does not support memory access across peer devices. The output parameter(s) are set 0
Contact Support for help.
Review the logic and adjust it manually.
DPCT1032
A different generator is used, you may need to adjust the code.
Intel® oneAPI Math Kernel Library
(
oneMKL
) RNG currently does not support the following engine types:
  • XORWOW generator
  • Sobol generator of 64-bit sequences
  • Scrambled Sobol generator
The engine types are migrated to the following supported generators, respectively:
  • Philox4x32x10 generator
  • Basic Sobol generator
  • Basic Sobol generator
You may need to adjust the code.
DPCT1033
The migrated code uses a basic Sobol generator. Initialize the
mkl::rng::sobol
generator with user-defined direction numbers to use it as Scrambled Sobol generator.
oneMKL
RNG currently does not support Scrambled Sobol generator by default.
Migrated code uses a basic Sobol generator.
Set user-defined direction numbers to the basic Sobol generator and use it as Scrambled Sobol generator.
See the Random Number Generators topic for more information.
DPCT1034
The migrated API does not return an error code. 0 is returned in the lambda. You may need to rewrite this code.
Typically, this happens because the API call in the original application returns an error code and then it is consumed by the program logic.
SYCL uses exceptions to report errors and does not return the error code.
The
DPC++
Compatibility Tool inserts a
return 0;
statement at the end of the lambda expression, if the return code is expected by the program logic and the new API does not return it. Review all such places in the code.
Similar to DPCT1003.
If in a
DPC++
application you:
  • Do not need the code that consumes the error code, remove the code and the
    return 0;
    statement.
  • Need the code that consumes the error code, try to replace it with an exception handling code and use your logic in an exception handler.
Similar to DPCT1003.
DPCT1035
All
DPC++
devices can be used by the host to submit tasks. You may need to adjust this code.
In
DPC++
, the host can freely submit tasks to all
DPC++
devices.
Review the code and adjust accordingly.
DPCT1036
The type
<type name>
was not migrated, because the migration depends on the
<API call>
.
The
Intel®
DPC++
Compatibility Tool
was unable to analyze and migrate the
<API call>
, so the
<type name>
type was not migrated.
See DPCT1028 for more details.
Rewrite this code manually by using a supported random number generator.
See the Random Number Generators topic for more information.
DPCT1037
Rewrite this code using
<library name>
with
DPC++
.
The
Intel®
DPC++
Compatibility Tool
does not support the migration of some library API calls.
Rewrite this code manually by using the
Intel® oneAPI Deep Neural Network Library
(
oneDNN
) or the
Intel® oneAPI Collective Communications Library
(
oneCCL
).
DPCT1038
When the kernel function name is used as a macro argument, the migration result may be incorrect. You need to verify the definition of the macro.
When a kernel function call is used in a macro definition and the function name is passed as a macro argument, the tool cannot determine the type of function parameters. This may lead to incorrectly generated
DPC++
code.
Review the kernel call inside the macro definition and adjust it manually, if needed.
DPCT1039
The generated code assumes that
<parameter name>
points to the global memory address space. If it points to a local memory address space, replace
<function name>
with
<function name>
.
The
Intel®
DPC++
Compatibility Tool
deduces whether the first parameter of an atomic function points to a global memory address space or a local memory space, using the last assignment's rvalue of the first parameter of the atomic function. If the last assignment is in an
if/while/do
while/for
statement, the deduction result may be incorrect. You need to verify the generated code to determine if the first parameter of the atomic function actually points to the local memory address space. If it does, then replace the atomic function name with an atomic function name that includes the template parameters, as pointed to in the warning message.
If the first parameter of an atomic function points to a local memory address space, replace the atomic function name with an atomic function name that includes the template parameters.
DPCT1040
Use
sycl::stream
instead of
printf
, if your code is used on the device.
If the
printf
statement is used on the host and the device in your original code, it does not change. To create output in DPC++,
sycl::stream
must be used on the device and
printf
must be used on the host.
  • If the
    printf
    statement is only used from the host, do not change your code.
  • If the
    printf
    statement is only used from the device, use
    sycl::stream
    instead of
    printf
    .
DPCT1041
SYCL uses exceptions to report errors, it does not use error codes. 0 is used instead of an error code in an if/while/do/for/switch/return statement. You may need to rewrite this code.
SYCL uses exceptions to report errors, it does not use error codes. The original code tries to get an error code, but SYCL does not require this functionality. Therefore, 0 is used instead of an error code in an if/while/do/for/switch/return statement.
You may need to rewrite this code.
DPCT1042
The size of the arguments passed to the SYCL kernel exceeds the minimum size limit (1024) for a non-custom SYCL device. You can get the hardware argument size limit by querying
info::device::max_parameter_size
. You may need to rewrite this code if the size of the arguments exceeds the hardware limit.
The size of the arguments passed to the SYCL kernel for non-custom SYCL device has a limit (see SYCL 1.2.1 standard, 4.6.4.2 Device information descriptors).
In cases where this warning occurs, you need to adjust the code manually to decrease the number of accessors or other arguments that are captured by the SYCL kernel lambda.
The following example shows how you can remove one accessor by merging two buffers with the same type.
Below is a problem example code where the arguments captured by the SYCL kernel lambda exceed the size limit:
... //declarations for device0 to device30 dpct::constant_memory<int, 1> device31(ARRAY_SIZE); dpct::constant_memory<int, 1> device32(ARRAY_SIZE); // kernel function declaration void kernel(sycl::nd_item<3> item_ct1, int *device0, ..., int *device32); void test_function() { ... //memcpy operations for device0 to device30 dpct::dpct_memcpy(device31.get_ptr(), host, ARRAY_SIZE * sizeof(int)); dpct::dpct_memcpy(device32.get_ptr(), host, ARRAY_SIZE * sizeof(int)); { dpct::get_default_queue().submit([&](sycl::handler &cgh) { ... //accessors for device0 to device30, captured by SYCL kernel lambda auto device31_acc_ct1 = device31.get_access(cgh); auto device32_acc_ct1 = device32.get_access(cgh); cgh.parallel_for( sycl::nd_range<3>(sycl::range<3>(1), sycl::range<3>(1)), [=](sycl::nd_item<3> item_ct1) { kernel(item_ct1, ... /*arguments device0 to device30*/ device31_acc_ct1.get_pointer(), device32_acc_ct1.get_pointer()); }); }); } }
Below is a solution example code where two buffers (device31 and device32) of the same type are merged into one (device31), and the number of accessors is reduced by one:
... //declarations for device0 to device30 dpct::constant_memory<int, 1> device32(ARRAY_SIZE + ARRAY_SIZE); // one buffer instead of two // kernel function declaration void kernel(sycl::nd_item<3> item_ct1, int *device0, ..., int *device32); void test_function() { ... //memcpy operations for device0 to device30 dpct::dpct_memcpy(device31.get_ptr(), host, ARRAY_SIZE * sizeof(int)); // memory copy with offset: dpct::dpct_memcpy(device31.get_ptr() + ARRAY_SIZE, host, ARRAY_SIZE * sizeof(int)); { dpct::get_default_queue().submit([&](sycl::handler &cgh) { ... //accessors for device0 to device30, captured by SYCL kernel lambda auto device31_acc_ct1 = device31.get_access(cgh); // only one accessor instead of 2 cgh.parallel_for( sycl::nd_range<3>(sycl::range<3>(1), sycl::range<3>(1)), [=](sycl::nd_item<3> item_ct1) { // last parameter is modified to use offset and the same accessor as previous parameter kernel(item_ct1, ... /*arguments device0 to device30*/ device31_acc_ct1.get_pointer(), device31_acc_ct1.get_pointer() + ARRAY_SIZE); }); }); } }
Review the code and adjust it.
DPCT1043
The version-related API is different in SYCL. An initial code was generated, but you need to adjust it.
The generated code uses
sycl::info::device::version
, which provides different information and uses a different type than the original code.
Adjust the generated code
DPCT1044
<BaseClass1>
was removed because
<BaseClass2>
has been deprecated in C++11. You may need to remove references to typedefs from
<BaseClass1>
in the class definition.
<BaseClass#>
is either
unary_function
or
binary_function
. The std:: equivalents for these classes are deprecated, and the use of these base classes is removed.
If any of the typedef identifiers are refenced in the class definition, they should be replaced with the original template arguments.
For example:
class C : thrust::unary_function<int,float> { argument_type arg_data; result_type result_data; }; should be rewritten as: class C { int arg_data; float result_data; };
DPCT1045
Migration is only supported for this API for general sparse matrix type. You may need to adjust
<original API name>
this code.
This warning appears if the used matrix type is not general, or cannot be determined.
If the matrix type used is:
  • Non-general: manually fix the code according to sparse-blas-routines.
  • General: ignore this warning.
DPCT1046
The
<original API name>
was not migrated because
<reason>
. You need to adjust the code. This may be due to one of the following:
  • Not all values of parameters could be evaluated in migration.
  • The combination of matrix data type and scalar type is unsupported.
Not all data type combinations are supported by
mkl::blas::gem()
. Use a supported data type to rewrite the code.
Please refer to the gemm topic of the Intel® oneAPI Math Kernel Library (oneMKL) - Data Parallel C++ Developer Reference for supported data types to fix the code manually.
DPCT1047
The meaning of
<parameter name>
in the
<API name>
is different from the
<API name>
. You may need to check the migrated code.
In cuBLAS* and cuSolver* getrf API, the LU factorization is done as
P*A=L*U
; in oneMKL API, it is done as
A=P*L*U
. The result of the matrix P may be different.
If the matrix P is only used in library API, ignore this warning. If P is used other ways, you may need to adjust the value of P.
DPCT1048
The original value
<macro name>
is not meaningful in the migrated code and was removed or replaced with 0. You may need to check the migrated code.
When API (calls, flags, etc.) are not meaningful in DPC++, they may be removed or replaced with 0, depending on how they are used.
Review the code and adjust it.
DPCT1049
The workgroup size passed to the SYCL kernel may exceed the limit. To get the device limit, query
info::device::max_work_group_size
. Adjust the workgroup size if needed.
The workgroup size passed to the SYCL kernel for SYCL device has a limit (see SYCL 1.2.1 standard, 4.6.4.2 Device information descriptors).
This warning appears if dimensions of the local range could not all be evaluated, or if the product of the dimensions of the local range is more than 256.
Query
info::device::max_work_group_size
to define the workgroup size limit for the device you use. If the workgroup size used in the code is below the limit, you can ignore this warning. Otherwise, you need to decrease the workgroup size.
DPCT1050
The template argument of the
<type>
could not be deduced. You need to update this code.
This warning is generated when the template argument could not be deduced by the Intel® DPC++ Compatibility Tool because the variable of this type was not used directly in the code. The Intel® DPC++ Compatibility Tool inserts "dpct_placeholder", instead of type, in such cases.
Replace the "dpct_placeholder" with the real argument.
DPCT1051
DPC++ does not support the device property that would be functionally compatible with
<property name>
. It was migrated to
<migrated token>
. You may need to rewrite the code.
Not all CUDA device properties currently have functionally compatible equivalents in DPC++. If such a property is detected and it is the integer type, it will be migrated to -1. If it is the Boolean type, it will be migrated to false or true depending on the context.
Review the logic and adjust it manually.
Contact Support for help.
DPCT1052
DPC++ does not support the member access for a volatile qualified vector type. The volatile qualifier was removed. You may need to rewrite the code.
The SYCL 1.2.1 specification does not provide a volatile qualified version of member access for vector type (see SYCL 1.2.1 standard, 4.10.2.1 Vec interface).
Review the code and adjust it.
Contact Support for help.
DPCT1053
Migration of device assembly code is not supported.
The assembly code was not replaced because it is hardware-specific. Replace this hardware-specific feature with the one offered by your hardware vendor.
Consult with your hardware vendor to find a replacement.
DPCT1054
The type of variable
<variable name>
is declared in device function with the name
<type>
. Adjust the code to make the
<type>
declaration visible at the accessor declaration point.
This warning will be emitted when the type of a __shared__ variable is declared in a device function.
In the migrated code:
- in the device function, the tool adds:
(a) uint8_t* input parameter
(b) name to the type, if unnamed, for example "type_ct1"
(c) a type cast from uint8_t* to the original type, after the original type declaration.
- in the call side, when defining local accessor, the first template argument will be set to uint8_t[sizeof(<original type, for example "type_ct1">)].
E.g.,
original code:
// header file template <typename T> __global__ void k(){ __shared__ union { T t; ... } a; ... } // source file void foo() { k<int><<<1,1>>>(); } migrated code: // header file template <typename T> void k(uint8_t *a_ct1) { union type_ct1 { T t; ... }; type_ct1* a = (type_ct1*)a_ct1; ... } // source file void foo() { dpct::get_default_queue().submit([&](sycl::handler &cgh) { /* DPCT1053:0: The type of variable a is declared in device function with the name type_ct1. Adjust the code to make the type_ct1 declaration visible at the accessor declaration point. */ sycl::accessor<uint8_t[sizeof(type_ct1)], 0, sycl::access::mode::read_write, sycl::access::target::local> a_ct1_acc_ct1(cgh); ... }); }
Move the type declaration so that it will be visible at the accessor declaration point or replace the "sizeof(<original type>)" with size in bytes needed for the original type.
DPCT1055
Vector types with size 1 are migrated to the corresponding fundamental types, which cannot be inherited. You need to rewrite the code.
The warning message is emitted when vector type with size 1 is inherited by a class or struct in the original code. Since the vector type with size 1 is migrated to the corresponding fundamental type in DPC++ and the fundamental type cannot be inherited, you need to rewrite the code.
You can declare a new field with the corresponding fundamental type, e.g. int for int1, in the class/struct and override the required operators.
For example, resulting code:
class MyClass : int{…} Fixed resulting code: class MyClass{ int x; MyClass operator+(const MyClass& y){...} MyClass operator=(const MyClass& y){...} ... }
DPCT1056
The Intel® DPC++ Compatibility Tool did not detect the variable
<variable name>
used in device code. If this variable is also used in device code, you need to rewrite the code.
If
__constant__
variable is only used in host code, "__constant__" attribute will be removed.
If this variable is only used in host code, ignore this warning.
If the variable is also used in device code, you need to migrate this code again with
--extra-arg="-xcuda"
passed to dpct to explicitly specify the code language.
DPCT1057
Variable
<variable name>
was used in host code and device code. The Intel® DPC++ Compatibility Tool updated
<variable name>
type to be used in SYCL device code and generated new
<host variable name>
to be used in host code. You need to update the host code manually to use the new
<host variable name>
.
If
__constant__variable
is used in both host code and device code (e.g., the variable is included in two compilation units and they are compiled by different compilers), it will be migrated to a
dpct::constant_memory
object and a new host variable
<host variable name>
.
You need to update the host code manually to use the new
<host variable name>
.
DPCT1058
<API Name>
is not migrated because it is not called in the code.
This warning is emitted when atomic API calls are used in macro definition and this macro is not called in the code.
If you want this macro to be migrated by Intel® DPC++ Compatibility Tool, please use the macro that needs migration in your code and pass the code to the tool again. Otherwise you can update the code manually or ignore the warning.
DPCT1059
SYCL only supports 4-channel image format. Adjust the code.
SYCL supports only 4-channel image format. The warning is emitted, when the tool generates code with unsupported image format, which corresponds to the original code. You can fix the resulting code by changing the image format. Note: suggested workaround may impact code performance.
Example:
Migrated code, which is using unsupported image format: dpct::image_wrapper<cl::sycl::uint2, 2> tex; //2-channel image is not supported void test_image(dpct::image_accessor_ext<cl::sycl::uint2, 2> acc){ cl::sycl::uint2 tex_data; tex_data = acc.read(0, 0); } int main(){ … dpct::get_default_queue().submit([&](cl::sycl::handler &cgh) { … auto acc = tex.get_access(cgh); auto smpl = tex.get_sampler(); … cgh.single_task<class dpct_single_kernel>([=]{ test_image(dpct::image_accessor_ext<cl::sycl::uint2, 2>(smpl, acc));} ); }); … } Manually fixed code: dpct::image_wrapper<cl::sycl::uint4, 2> tex; void test_image(dpct::image_accessor_ext<cl::sycl::uint4, 2> acc){ cl::sycl::uint4 tex_data; tex_data = acc.read(0, 0); } int main(){ … dpct::get_default_queue().submit([&](cl::sycl::handler &cgh) { … auto acc = tex.get_access(cgh); auto smpl = tex.get_sampler(); … cgh.single_task<class dpct_single_kernel>([=]{ test_image(dpct::image_accessor_ext<cl::sycl::uint4, 2>(smpl, acc));} ); }); … }
You may need to rewrite this code.
DPCT1060
SYCL range can only be a 1D, 2D, or 3D vector. Adjust the code.
This warning is emitted when the number of dimensions of memory in the original code exceeds 3. Since SYCL range supports only 1, 2 or 3 dimensions, the resulting code is not SYCL-compliant.
To fix the resulting code you can use the low-dimensional arrays to simulate high-dimensional arrays.
The following fix example demonstrates how to use a 3D array to simulate a 4D array.
Resulting Code:
dpct::constant_memory<int, 4> array(dimX, dimY, dimZ, dimW); void kernel(sycl::id<1> idx, dpct::accessor<int, dpct::constant, 4> const_array) { … … = const_array[x][y][z][w]; … } Manually fixed code: dpct::constant_memory<int, 3> array(dimX, dimY, dimZ * dimW); void kernel(sycl::id<1> idx, dpct::accessor<int, dpct::constant, 3> const_array) { … … = const_array[x][y][w * dimZ + z]; … }
You may need to rewrite this code.
DPCT1061
Call to
<macro name>
macro was removed, because it only contains code, which is unnecessary in DPC++.
When a macro definition contains only code, which will be removed during migration, DPCT removes the usage of the macro as well.
Review the code and remove the macro definition if it is not needed.
DPCT1062
SYCL Image doesn't support normalized read mode.
This warning is emitted when
cudaReadModeNormalizedFloat
is used as the third argument of texture in the original code. Since SYCL Image doesn't support normalized read mode,
cudaReadModeNormalizedFloat
will be ignored during migration.
It may cause errors in the resulting code, for example, redefinition of overloaded functions, if the overloaded functions are differentiated based on the texture type in the original code.
Original code: __device__ void foo(const texture<char, 2, cudaReadModeNormalizedFloat> tex); __device__ void foo (const texture<char, 2, cudaReadModeElementType> tex); Resulting code: void foo (dpct::image_accessor_ext<char, 2> tex); void foo (dpct::image_accessor_ext<char, 2> tex);
Review the code and update as needed.
DPCT1063
Advice parameter is device-defined. You may need to adjust it.
The advice value is device-defined advice for the specified memory allocation. A value of 0 reverts the advice to the default behavior. Replace the value with the one required by your device.
Consult with your hardware vendor to find a replacement.
DPCT1064
Migrated
<function name>
call is used in a macro definition and is not valid for all macro uses. Adjust the code.
The warning is generated when the function call is used in a macro definition and needs to be migrated differently, depending on how the macro is called. The Intel® DPC++ Compatibility Tool generates code that is valid for one of the calls of the macro, but may not be valid for all calls of this macro in the code.
For example: The function
pow
can be migrated to
sycl::pow<double>
or
sycl::pown
, depending on the types of parameters passed through macro arguments.
#define POW(B, E) pow(B, E) POW(2.5, 3.1); //should be migrated to sycl::pow<double>(2.5, 3.1) POW(2.5, 3); //should be migrated to sycl::pown((float)2.5, 3)
Declare new macros for different use cases of the macro call in the resulting code.

Product and Performance Information

1

Performance varies by use, configuration and other factors. Learn more at www.Intel.com/PerformanceIndex.