Skip to content

Commit a071554

Browse files
jinge90jinge90
authored
[SYCL] Add String utils in libdevice (#19356)
Signed-off-by: jinge90 <[email protected]>
1 parent cdb4cab commit a071554

File tree

4 files changed

+273
-30
lines changed

4 files changed

+273
-30
lines changed

‎libdevice/crt_wrapper.cpp

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,59 @@ int memcmp(const void *s1, const void *s2, size_t n) {
3232
return __devicelib_memcmp(s1, s2, n);
3333
}
3434

35+
// We align with libc string functions, no checking for null ptr.
36+
DEVICE_EXTERN_C_INLINE
37+
size_t strlen(const char *s) {
38+
const char *t;
39+
for (t = s; *t; ++t)
40+
;
41+
return t - s;
42+
}
43+
44+
DEVICE_EXTERN_C_INLINE
45+
char *strcpy(char *dest, const char *src) {
46+
char *t = dest;
47+
for (; (*dest = *src) != '\0'; ++dest, ++src)
48+
;
49+
return t;
50+
}
51+
52+
DEVICE_EXTERN_C_INLINE
53+
char *strncpy(char *dest, const char *src, size_t n) {
54+
size_t i;
55+
for (i = 0; i < n && (src[i] != '\0'); ++i)
56+
dest[i] = src[i];
57+
for (; i < n; ++i)
58+
dest[i] = '\0';
59+
return dest;
60+
}
61+
62+
DEVICE_EXTERN_C_INLINE
63+
int strcmp(const char *s1, const char *s2) {
64+
while (*s1 == *s2) {
65+
if (*s1 == '\0')
66+
return 0;
67+
++s1;
68+
++s2;
69+
}
70+
71+
return *reinterpret_cast<const unsigned char *>(s1) -
72+
*reinterpret_cast<const unsigned char *>(s2);
73+
}
74+
75+
DEVICE_EXTERN_C_INLINE
76+
int strncmp(const char *s1, const char *s2, size_t n) {
77+
78+
size_t idx = 0;
79+
while ((idx < n) && (s1[idx] != '\0') && (s1[idx] == s2[idx]))
80+
idx++;
81+
82+
if (idx == n)
83+
return 0;
84+
85+
return s1[idx] - s2[idx];
86+
}
87+
3588
// This simple rand is for ease of use only, the implementation aligns with
3689
// LLVM libc rand which is based on xorshift64star pseudo random number
3790
// generator. If work item number <= 1024, each work item has its own internal

‎sycl/include/sycl/builtins.hpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,13 @@ extern __DPCPP_SYCL_EXTERNAL_LIBC void *memcpy(void *dest, const void *src,
2020
extern __DPCPP_SYCL_EXTERNAL_LIBC void *memset(void *dest, int c, size_t n);
2121
extern __DPCPP_SYCL_EXTERNAL_LIBC int memcmp(const void *s1, const void *s2,
2222
size_t n);
23+
extern __DPCPP_SYCL_EXTERNAL_LIBC size_t strlen(const char *s);
24+
extern __DPCPP_SYCL_EXTERNAL_LIBC char *strcpy(char *dest, const char *src);
25+
extern __DPCPP_SYCL_EXTERNAL_LIBC char *strncpy(char *dest, const char *src,
26+
size_t n);
27+
extern __DPCPP_SYCL_EXTERNAL_LIBC int strcmp(const char *s1, const char *s2);
28+
extern __DPCPP_SYCL_EXTERNAL_LIBC int strncmp(const char *s1, const char *s2,
29+
size_t n);
2330
extern __DPCPP_SYCL_EXTERNAL_LIBC int rand();
2431
extern __DPCPP_SYCL_EXTERNAL_LIBC void srand(unsigned int seed);
2532
extern __DPCPP_SYCL_EXTERNAL long long int __imf_llmax(long long int x,

‎sycl/test-e2e/Basic/code_location_e2e.cpp

Lines changed: 4 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -22,37 +22,12 @@
2222
2323
*/
2424

25+
#include <cstring>
2526
#include <sycl/detail/core.hpp>
2627

2728
#include <sycl/stream.hpp>
2829
using namespace sycl;
2930

30-
// llvm/sycl/doc/design/DeviceLibExtensions.rst
31-
// Our devicelib support for <cstring> only includes three memory
32-
// operations, none of the string ones. So we need to provide
33-
// our own string comparison for kernel calls.
34-
bool stringsAreSameP(const char *a, const char *b) {
35-
// If both are nullptr, then they are the same,
36-
if ((a == nullptr) && (b == nullptr))
37-
return true;
38-
// but if only one, they are not.
39-
if ((a == nullptr) || (b == nullptr))
40-
return false;
41-
42-
int index = 0;
43-
while (true) {
44-
if (a[index] != b[index]) {
45-
return false;
46-
}
47-
if (a[index] == '\0') {
48-
return true;
49-
} // If we are on this line we know a[i]==b[i].
50-
index++;
51-
}
52-
// We will never arrive here.
53-
return true;
54-
}
55-
5631
template <typename OS> void report(OS &out, detail::code_location code_loc) {
5732
out << "function {line:col} => " << code_loc.functionName() << " {"
5833
<< code_loc.lineNumber() << ":" << code_loc.columnNumber() << "}"
@@ -73,9 +48,8 @@ void test(OS &out, detail::code_location &code_loc, const char *fileName,
7348
// functionName
7449
auto funcNameStr = code_loc.functionName();
7550
auto fNameResult =
76-
((funcNameStr != nullptr) && stringsAreSameP(funcNameStr, funcName))
77-
? "OK"
78-
: "WRONG";
51+
((funcNameStr != nullptr) && !strcmp(funcNameStr, funcName)) ? "OK"
52+
: "WRONG";
7953
out << "code_location.functionName: " << fNameResult << "\n";
8054

8155
// lineNumber
@@ -95,7 +69,7 @@ void test(OS &out, detail::code_location &code_loc, const char *fileName,
9569
? "OK"
9670
: "WRONG - fileName should not be present when NDEBUG defined";
9771
#else
98-
auto fileNameResult = stringsAreSameP(fileName, fileNameStr) ? "OK" : "WRONG";
72+
auto fileNameResult = (strcmp(fileName, fileNameStr) == 0) ? "OK" : "WRONG";
9973
#endif
10074
out << "code_location.fileName: " << fileNameResult << "\n";
10175
}

‎sycl/test-e2e/DeviceLib/string_test.cpp

Lines changed: 209 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,200 @@ bool kernel_test_memcpy(sycl::queue &deviceQueue) {
5252
return success;
5353
}
5454

55+
class KernelTestStrcpy;
56+
bool kernel_test_strcpy(sycl::queue &deviceQueue) {
57+
bool success = true;
58+
char src[20] = "abcdefg012345xyzvvv";
59+
char dst[5][20];
60+
using CStr = char[20];
61+
{
62+
sycl::buffer<CStr, 1> dst_buffer(dst, sycl::range<1>(5));
63+
deviceQueue.submit([&](sycl::handler &cgh) {
64+
auto dst_acc = dst_buffer.get_access<sycl::access::mode::write>(cgh);
65+
cgh.single_task<class KernelTestStrcpy>([=]() {
66+
char str[20] = "abcdefg012345xyzvvv";
67+
strcpy(dst_acc[0], str);
68+
str[17] = '\0';
69+
strcpy(dst_acc[1], str);
70+
str[12] = '\0';
71+
strcpy(dst_acc[2], str);
72+
str[8] = '\0';
73+
strcpy(dst_acc[3], str);
74+
str[0] = '\0';
75+
strcpy(dst_acc[4], str);
76+
});
77+
});
78+
}
79+
80+
if (strcmp(src, dst[0]) != 0)
81+
success = false;
82+
src[17] = '\0';
83+
if (strcmp(src, dst[1]) != 0)
84+
success = false;
85+
src[12] = '\0';
86+
if (strcmp(src, dst[2]) != 0)
87+
success = false;
88+
src[8] = '\0';
89+
if (strcmp(src, dst[3]) != 0)
90+
success = false;
91+
src[0] = '\0';
92+
if (strcmp(src, dst[4]) != 0)
93+
success = false;
94+
return success;
95+
}
96+
97+
class KernelTestStrncpy;
98+
bool kernel_test_strncpy(sycl::queue &deviceQueue) {
99+
char src[20] = "abcdefg012345xyzvvv";
100+
char dst[3][20];
101+
memset(reinterpret_cast<char *>(dst), 'A', 60);
102+
typedef char CStr[20];
103+
{
104+
sycl::buffer<CStr, 1> dst_buffer(dst, sycl::range<1>(3));
105+
deviceQueue.submit([&](sycl::handler &cgh) {
106+
auto dst_acc = dst_buffer.get_access<sycl::access::mode::write>(cgh);
107+
cgh.single_task<class KernelTestStrncpy>([=]() {
108+
char str[20] = "abcdefg012345xyzvvv";
109+
strncpy(dst_acc[0], str, 19);
110+
strncpy(dst_acc[1], str, 20);
111+
str[7] = '\0';
112+
strncpy(dst_acc[2], str, 11);
113+
});
114+
});
115+
}
116+
117+
size_t idx;
118+
for (idx = 0; idx < 19; ++idx) {
119+
if (dst[0][idx] != src[idx])
120+
return false;
121+
}
122+
123+
if (dst[0][19] != 'A')
124+
return false;
125+
126+
for (idx = 0; idx < 20; ++idx) {
127+
if (dst[1][idx] != src[idx])
128+
return false;
129+
}
130+
131+
for (idx = 0; idx < 7; ++idx) {
132+
if (dst[2][idx] != src[idx])
133+
return false;
134+
}
135+
for (idx = 7; idx < 11; ++idx) {
136+
if (dst[2][idx] != '\0') {
137+
return false;
138+
}
139+
}
140+
141+
if (dst[2][11] != 'A')
142+
return false;
143+
144+
return true;
145+
}
146+
147+
class KernelTestStrcmp;
148+
bool kernel_test_strcmp(sycl::queue &deviceQueue) {
149+
int res[5];
150+
{
151+
sycl::buffer<int, 1> res_buffer(res, sycl::range<1>(5));
152+
deviceQueue.submit([&](sycl::handler &cgh) {
153+
auto res_acc = res_buffer.get_access<sycl::access::mode::write>(cgh);
154+
cgh.single_task<class KernelTestStrcmp>([=]() {
155+
char str1[20] = "abcdefg012";
156+
char str2[20] = "abcd";
157+
char str3[20] = "124ddf";
158+
char str4[20] = "abcdefg015";
159+
res_acc[0] = strcmp(str1, str1);
160+
res_acc[1] = strcmp(str1, str2);
161+
res_acc[2] = strcmp(str3, str1);
162+
res_acc[3] = strcmp(str4, str1);
163+
res_acc[4] = strcmp(str1, str4);
164+
});
165+
});
166+
}
167+
168+
if ((res[0] != 0) || (res[1] <= 0) || (res[2] >= 0) || (res[3] <= 0) ||
169+
(res[4] >= 0))
170+
return false;
171+
return true;
172+
}
173+
174+
class KernelTestStrncmp;
175+
bool kernel_test_strncmp(sycl::queue &deviceQueue) {
176+
int res[10];
177+
{
178+
sycl::buffer<int, 1> res_buffer(res, sycl::range<1>(10));
179+
deviceQueue.submit([&](sycl::handler &cgh) {
180+
auto res_acc = res_buffer.get_access<sycl::access::mode::write>(cgh);
181+
cgh.single_task<class KernelTestStrncmp>([=]() {
182+
char str1[20] = "abcdefg012";
183+
char str2[20] = "abcd";
184+
char str3[20] = "124ddf";
185+
char str4[20] = "abcdefg015";
186+
char str5[20] = "";
187+
char str6[20] = "abcdEFG";
188+
char str7[20] = "abcdefg";
189+
str6[3] = str7[3] = '\0';
190+
res_acc[0] = strncmp(str1, str1, 12);
191+
res_acc[1] = strncmp(str1, str2, 4);
192+
res_acc[2] = strncmp(str3, str1, 0);
193+
res_acc[3] = strncmp(str4, str1, 9);
194+
res_acc[4] = strncmp(str1, str4, 10);
195+
res_acc[5] = strncmp(str3, str3, 20);
196+
res_acc[6] = strncmp(str2, str3, 6);
197+
res_acc[7] = strncmp(str5, str5, 12);
198+
res_acc[8] = strncmp(str6, str7, 7);
199+
str2[0] = str3[0] = '\0';
200+
res_acc[9] = strncmp(str2, str3, 4);
201+
});
202+
});
203+
}
204+
205+
if ((res[0] != 0) || (res[1] != 0) || (res[2] != 0) || (res[3] != 0) ||
206+
(res[4] >= 0) || (res[5] != 0) || (res[6] <= 0) || (res[7] != 0) ||
207+
(res[8] != 0) || (res[9] != 0))
208+
return false;
209+
return true;
210+
}
211+
212+
class KernelTestStrlen;
213+
bool kernel_test_strlen(sycl::queue &deviceQueue) {
214+
bool success = true;
215+
char src[20] = "abcdefg012345xyzvvv";
216+
size_t len[5] = {0};
217+
{
218+
sycl::buffer<char, 1> buffer1(src, sycl::range<1>(20));
219+
sycl::buffer<size_t, 1> buffer2(len, sycl::range<1>(5));
220+
deviceQueue.submit([&](sycl::handler &cgh) {
221+
auto len_acc = buffer2.get_access<sycl::access::mode::write>(cgh);
222+
auto src_acc = buffer1.get_access<sycl::access::mode::read_write>(cgh);
223+
cgh.single_task<class KernelTestStrlen>([=]() {
224+
len_acc[0] =
225+
strlen(src_acc.get_multi_ptr<sycl::access::decorated::no>().get());
226+
src_acc[17] = '\0';
227+
len_acc[1] =
228+
strlen(src_acc.get_multi_ptr<sycl::access::decorated::no>().get());
229+
src_acc[12] = '\0';
230+
len_acc[2] =
231+
strlen(src_acc.get_multi_ptr<sycl::access::decorated::no>().get());
232+
src_acc[7] = '\0';
233+
len_acc[3] =
234+
strlen(src_acc.get_multi_ptr<sycl::access::decorated::no>().get());
235+
src_acc[0] = '\0';
236+
len_acc[4] =
237+
strlen(src_acc.get_multi_ptr<sycl::access::decorated::no>().get());
238+
});
239+
});
240+
}
241+
242+
if ((len[0] != 19) || (len[1] != 17) || (len[2] != 12) || (len[3] != 7) ||
243+
(len[4] != 0))
244+
success = false;
245+
246+
return success;
247+
}
248+
55249
class KernelTestMemcpyInit;
56250
class KernelTestMemcpyUSM0;
57251
class KernelTestMemcpyUSM1;
@@ -489,6 +683,21 @@ int main() {
489683

490684
success = kernel_test_memcpy_addr_space(deviceQueue);
491685
assert(((void)"memcpy test with address space failed!", success));
686+
687+
success = kernel_test_strlen(deviceQueue);
688+
assert(((void)"strlen test failed!", success));
689+
690+
success = kernel_test_strcpy(deviceQueue);
691+
assert(((void)"strcpy test failed!", success));
692+
693+
success = kernel_test_strncpy(deviceQueue);
694+
assert(((void)"strncpy test failed!", success));
695+
696+
success = kernel_test_strcmp(deviceQueue);
697+
assert(((void)"strcmp test failed!", success));
698+
699+
success = kernel_test_strncmp(deviceQueue);
700+
assert(((void)"strncmp test failed!", success));
492701
std::cout << "passed!" << std::endl;
493702
return 0;
494703
}

0 commit comments

Comments
 (0)