string.cuh

/**
 * @file on_device/c_standard_library/string.cuh
 * @brief CUDA device-side versions of the C standard library string operations.
 *
 * @note The implementations are not intended to be particularly speedy - merely functional.
 * In particular - you are advised to only use this when each thread is working on its own
 * string.
 *
 * @note unimplemented functions: strcoll, strxfrm (these require a locale); strerror (errno
 * errors are not generated by device-side code); strtok (not reentrant)
 *
 * @note The implementations are not actually C functions; I've put them in C++ namespaces.
 * This choice could be changed if users demand it.
 */
#pragma once
#ifndef CUDAT_KAT_ON_DEVICE_C_STANDARD_LIBRARY_EQUIVALENTS_STRING_H_
#define CUDAT_KAT_ON_DEVICE_C_STANDARD_LIBRARY_EQUIVALENTS_STRING_H_

#include <kat/on_device/common.cuh>


///@cond
#include <kat/detail/execution_space_specifiers.hpp>
///@endcond

namespace kat {

namespace c_std_lib {

using size_t = std::size_t;

/**
 * @brief Lexicographically compares two nul-terminated strings.
 *
 * @note comparison is not locale-specific
 *
 * @return a negative value if lhs > rhs, zero if lhs == rhs and a positive
 * value if lhs < rhs.
 */
inline __device__
int strcmp(const char* lhs, const char* rhs)
{
    while((*lhs != '\0') and (*rhs == *lhs)) {
		lhs++, rhs++;
    }
    return ((unsigned char) *lhs) - ((unsigned char) *rhs);
}

/**
 * @brief Lexicographically compares two nul-terminated strings - but no farther than their n'th character.
 *
 * @note comparison is not locale-specific.
 *
 * @return a negative value if lhs > rhs, zero if lhs == rhs and a positive
 * value if lhs < rhs - but with the comparisons ignoring everything after the n'th character of each
 * of the strings, so - if they differ after then n'th character, zero is returned.
 */
inline __device__
int strncmp(const char* lhs, const char* rhs, size_t n)
{
	while (n && *lhs && (*lhs == *rhs)) {
		++lhs;
		++rhs;
		--n;
	}
	if (n == 0) { return 0; }
	return (*(unsigned char *) lhs - *(unsigned char *) rhs);
}

/**
 * @brief Lexicographically compare two equal-length sequences of bytes.
 *
 * @note similar to strncmp, except that the comparison does not stop with 0-valued bytes.
 *
 * @param lhs
 * @param rhs
 * @param n The length of each sequence of bytes
 * @return a negative value if lhs > rhs, zero if lhs == rhs and a positive
 * value if lhs < rhs
 */
inline __device__
int memcmp(const void* lhs, const void* rhs, size_t n)
{
	auto lhs_as_chars = static_cast<const char*>(lhs);
	auto rhs_as_chars = static_cast<const char*>(rhs);
	for(size_t i = 0; i < n; i++) {
		if (lhs_as_chars[i] > rhs_as_chars[i]) { return  1; }
		if (lhs_as_chars[i] < rhs_as_chars[i]) { return -1; }
	}
	return 0;
}

/**
 * @brief Copy a nul-terminated string from one location to another
 *
 * @return the target location (which, after execution, should contain
 * a copy of the string --- when accessed from the calling thread or
 * after appropriate synchronization).
 */
inline __device__
char* strcpy(char *dst, const char *src)
{
	auto ret = dst;
	while (*src != '\0') { *(dst++) = *(src++); }
	*dst = *src;
	return ret;
}

/**
 * @brief Copy a nul-terminated string from one location to another - but
 * always write no more and no less than a fixed number of characters.
 *
 * @param dst copy destination
 * @param src string to copy from
 * @param n the number of characters to write at the destination
 * @return the target location (which, after execution, should contain
 * a copy of the string or the first @p n characters thereof --- when
 * accessed from the calling thread or after appropriate synchronization).
 */
inline __device__
char* strncpy(char *dst, const char *src, size_t n)
{
	size_t i = 0;
	auto ret = dst;
	for(; i < n && *src != '\0'; i++, src++, dst++) {
		*dst = *src;
	}
	for(; i < n; i++, dst++) {
		*dst = '\0';
	}
	return ret;
}

/**
 * @brief determines the length of a nul-terminated string
 *
 * @return The length of the string, i.e. the number of characters
 * before the terminating '\0' (nul) character.
 */
inline __device__
std::size_t strlen(const char *s)
{
	const char* p = s;
	while(*p != '\0') { p++; }
	return p - s;
}

/**
 * @brief concatenates one string after another, in-place.
 *
 * @param dest The target location for both string (after concatenation)
 * @param src The string to be concatenated.
 * @return A pointer to @p dest, after it has gotten a copy of @p src
 * concatenated at its end.
 */
inline __device__
char *strcat(char *dest, const char *src)
{
	strcpy(dest + strlen(dest), src);
	return dest;
}

/**
 * @brief concatenates one string after another, in-place - but appending
 * no more than a fixed number of characters to the string at the destination.
 *
 * @param dest The target location for both string (after concatenation)
 * @param src The string to be concatenated.
 * @param n Maximum number of characters used from the source string
 * @return A pointer to @p dest, after it has gotten a copy of @p src
 * concatenated at its end - or the first @p n characters of @p src
 * followed with a terminating '\0' (nul).
 */
inline __device__
char *strncat(char *dest, const char *src, size_t n)
{
	strncpy(dest + strlen(dest), src, n);
	return dest;
}

/**
 * @brief copies a stretch of memory from one location to another
 *
 * @note the two regions - @p size bytes at @p source, @p size
 * bytes at @p destination - must be disjoint
 *
 * @param destination beginning of the target region
 * @param source beginning of the source region
 * @param size the length of the memory region to copy
 */
inline __device__
void* memcpy(
	void*        __restrict__  destination,
	const void*  __restrict__  source,
	size_t                     size)
{
	return ::memcpy(destination, source, size);
}

/**
 * @brief set a sequence of bytes in memory to a fixed value
 *
 * @param destination beginning of the sequence of bytes to set
 * @param c value to which to set the bytes; must be in the range
 * of 0 .. (1 << CHAR_BIT) - 1 .
 * @param size number of bytes to set to value @p c.
 */
inline __device__
void* memset(void* destination, int c, size_t size)
{
	::memset(destination, c, size);
	return destination;
}

/**
 * @brief locate a byte with a specified value in a sequence of
 * bytes in memory.
 * @param s beginning of the sequence of bytes at which to begin the search
 * @param c value to search for; must be in the range
 * of 0 .. (1 << CHAR_BIT) - 1 .
 * @param n number of bytes to search for the desired value
 * @return address of the first byte in the sequence with the value @p c,
 * within the sequence starting at @p s, or nullptr if none of the first
 * @p n bytes have that value.
 */
inline __device__
void *memchr(const void *s, int c, size_t n)
{
	auto s_end = ((const char *) s) + n;
	for(const char* p = (const char *) s; p < s_end; p++) {
		if (*p == c) { return (void *) p; }
	}
	return nullptr;
}

/**
 * @brief Search for a character within a nul-terminated string.
 *
 * @param s The string to search
 * @param c A character value to search for
 * @return address of the first character with the value @p c
 * within string @p s, or nullptr if no character of @p s equals @p c .
 */
inline __device__
char *strchr(const char *s, int c)
{
	const char* p = s;
	do {
		if (*p == static_cast<char>(c)) {
			return const_cast<char*>(p);
		}
	} while(*(p++) != '\0');
	return nullptr;
}

/**
 * @brief same as @ref std::strchr , except that the search begins
 * at the end of the string
 *
 * @note If @p c is '\0', it _will_ match the nul character
 * at the end of the string.
 *
 */
inline __device__
char *strrchr(const char *s, int c)
{
	const char* last = nullptr;
	const char* p = s;
	do {
		if (*p == c) { last = p; }
	} while(*(p++) != '\0');
	return const_cast<char*>(last);
}

// Naive implementation!
inline __device__
char *strpbrk(const char *s, const char *accept)
{
	for(const char* p = s; *p != '\0'; *p++) {
		if (strchr(accept, *p)) { return const_cast<char*>(p); }
	}
	return nullptr;
}

// Naive implementation!
inline __device__
size_t strspn(const char *s, const char *accept)
{
	const char* p = s;
	while(*p != '\0' && strchr(accept, *p)) { p++; }
	return p - s;
}

// Naive implementation!
inline __device__
size_t strcspn(const char *s, const char *reject)
{
	const char* p = s;
	while(*p != '\0' && !strchr(reject, *p)) { p++; }
	return p - s;
}

// Naive O(|haystack| * |needle|) implementation!
inline __device__
char *strstr(const char *haystack, const char *needle)
{
	auto match_prefix = [](const char* s, const char* prefix) {
	    while((*prefix != '\0') and (*s == *prefix)) {
			s++, prefix++;
	    }
	    return (*prefix == '\0');
	};

	do {
		if (match_prefix(haystack, needle)) {
			return const_cast<char *>(haystack);
		}
	} while(*(haystack++) != '\0');
	return (*needle == '\0') ? const_cast<char *>(haystack) : nullptr;
}

// Naive O(|haystack| * |needle|) implementation!
inline __device__
char *strrstr(const char *haystack, const char *needle)
{
	auto match_prefix = [](const char* s, const char* prefix) {
	    while((*prefix != '\0') and (*s == *prefix)) {
			s++, prefix++;
	    }
	    return (*prefix == '\0');
	};

	const char* last_match = nullptr;
	do {
		if (match_prefix(haystack, needle)) {
			last_match = haystack;
		}
	} while(*(haystack++) != '\0');
	return const_cast<char *>((*needle == '\0') ? haystack : last_match);
}



} // namespace c_std_lib

} // namespace kat


#include <kat/detail/execution_space_specifiers.hpp>

#endif // CUDAT_KAT_ON_DEVICE_C_STANDARD_LIBRARY_EQUIVALENTS_STRING_H_