Visible to Intel only — GUID: GUID-18728437-120B-4FE5-9344-3E5340F1BB08
DPCT1000
DPCT1001
DPCT1002
DPCT1003
DPCT1004
DPCT1005
DPCT1006
DPCT1007
DPCT1008
DPCT1009
DPCT1010
DPCT1011
DPCT1012
DPCT1013
DPCT1014
DPCT1015
DPCT1016
DPCT1017
DPCT1018
DPCT1019
DPCT1020
DPCT1021
DPCT1022
DPCT1023
DPCT1024
DPCT1025
DPCT1026
DPCT1027
DPCT1028
DPCT1029
DPCT1030
DPCT1031
DPCT1032
DPCT1033
DPCT1034
DPCT1035
DPCT1036
DPCT1037
DPCT1038
DPCT1039
DPCT1040
DPCT1041
DPCT1042
DPCT1043
DPCT1044
DPCT1045
DPCT1046
DPCT1047
DPCT1048
DPCT1049
DPCT1050
DPCT1051
DPCT1052
DPCT1053
DPCT1054
DPCT1055
DPCT1056
DPCT1057
DPCT1058
DPCT1059
DPCT1060
DPCT1061
DPCT1062
DPCT1063
DPCT1064
DPCT1065
DPCT1066
DPCT1067
DPCT1068
DPCT1069
DPCT1070
DPCT1071
DPCT1072
DPCT1073
DPCT1074
DPCT1075
DPCT1076
DPCT1077
DPCT1078
DPCT1079
DPCT1080
DPCT1081
DPCT1082
DPCT1083
DPCT1084
DPCT1085
DPCT1086
DPCT1087
DPCT1088
DPCT1089
DPCT1090
DPCT1091
DPCT1092
DPCT1093
DPCT1094
DPCT1095
DPCT1096
Message
Detailed Help
Suggestions to Fix
DPCT1097
DPCT1098
DPCT1099
DPCT1100
Visible to Intel only — GUID: GUID-18728437-120B-4FE5-9344-3E5340F1BB08
DPCT1096
Message
The right-most dimension of the work-group used in the SYCL kernel that calls this function may be less than <value of kernel sub-group size attribute>. The function <help function name> may return an unexpected result on the CPU device. Modify the size of the work-group to ensure that the value of the right-most dimension is a multiple of <value of kernel sub-group size attribute>.
Detailed Help
The dpct::select_from_sub_group, dpct::shift_sub_group_left, dpct::shift_sub_group_right, and dpct::permute_sub_group_by_xor functions may return unexpected results when run on a CPU device with an OpenCL™ backend, if the right-most dimension value of the work-group used in the SYCL* kernel that calls these functions is less than the value of the kernel sub-group size attribute. The real sub-group size may not be the value specified by the kernel sub-group size attribute and could cause the helper function to return unexpected results on a CPU device.
Adjust the code by modifying the size of the work-group to ensure that the value of the right-most dimension is a multiple of the kernel sub-group size attribute.
For example, this original CUDA* code:
__global__ void kernel() {
...
value = __shfl_down(x, delta);
...
}
int main() {
...
auto GridSize = dim3(2);
auto BlockSize = dim3(8, 8, 1);
kernel<<<GridSize, BlockSize>>>();
...
}
results in the following migrated code:
void kernel(sycl::nd_item<3> item) {
...
value = dpct::shift_sub_group_left(item.get_sub_group(), x, delta); // May return unexpected result on CPU
...
}
int main() {
...
auto GridSize = sycl::range<3>(1, 1, 2);
auto BlockSize = sycl::range<3>(1, 8, 8); // Problem: value of the right-most dimension 8 is less than the kernel sub group size attribute 32.
queue.parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item)[[intel::reqd_sub_group_size(32)]] {
kernel(item);
});
...
}
which is rewritten to:
void kernel(sycl::nd_item<3> item) {
...
value = dpct::shift_sub_group_left(item.get_sub_group(), x, delta);
...
}
int main() {
auto GridSize = sycl::range<3>(1, 1, 2);
auto BlockSize = sycl::range<3>(1, 2, 32); // Fix: modified work group size to make the right-most dimension to be multiple of the kernel sub group size attribute value, which is 32.
queue.parallel_for(sycl::nd_range<3>(GridSize * BlockSize, BlockSize), [=](sycl::nd_item<3> item)[[intel::reqd_sub_group_size(32)]] {
kernel(item);
});
...
}
Suggestions to Fix
If the program needs to execute on a CPU device, you may need to adjust the code.