bencudastr.cu

GPU实现的MapReduce framework,对于学习并行编程和cuda平台的编程方面有着极好的参考价值

#ifndef __BEN_CUDA_STR_CU__
#define __BEN_CUDA_STR_CU__

__device__ char *StrCpy(char *dest, const char *src, int size)
{
	char4 *dest1 = (char4*)dest;
	const char4 *src1 = (char4*)src;

	int numChar4=size>>2;
	int i=0;
	for(i=0;i<numChar4;i++)
	{
		dest1[i]=src1[i];
	}
	for(i=(numChar4<<2);i<size;i++)
	{
		dest[i]=src[i];
	}
	return dest; 
}

__device__ size_t StrLen(char *str)
{
	char *p = NULL;
	for (p = str; *p != '\0'; p++);
	return (size_t)(p - str - 1);
}

__device__ char *StrCat(char *dest, int dest_size, char *src, int src_size)
{
	return StrCpy(dest+dest_size, src, src_size);
}

__device__ int StrCmp(const char *str1, int size1, const char *str2, int size2)
{
	int minSize = (size1 < size2) ? size1:size2;
	
	char4 *char4_str1 = (char4*)str1;
	char4 *char4_str2 = (char4*)str2;
	
	int numChar4=minSize>>2;
	int i=0;
	char4 curChar4_1;
	char4 curChar4_2;
	for(i=0;i<numChar4;i++)
	{
		curChar4_1=char4_str1[i];
		curChar4_2=char4_str2[i];
		if(curChar4_1.x>curChar4_2.x) return 1;
		else if(curChar4_1.x<curChar4_2.x) return -1;

		if(curChar4_1.y>curChar4_2.y) return 1;
		else if(curChar4_1.y<curChar4_2.y) return -1;

		if(curChar4_1.z>curChar4_2.z) return 1;
		else if(curChar4_1.z<curChar4_2.z) return -1;

		if(curChar4_1.w>curChar4_2.w) return 1;
		else if(curChar4_1.w<curChar4_2.w) return -1;
	}
	char c1;
	char c2;
	for(i=(numChar4<<2);i<minSize;i++)
	{
		c1=str1[i]; c2=str2[i];
		if(c1>c2) return 1;
		else if(c1<c2) return -1;
	}
	if(size1<size2) return -1;
	else if(size1>size2) return 1;
	else return 0;
	
}

__device__ char *StrChr(char *str, int c)
{
	char *p = str;
	while (*p != '\0')
	{
		if (*p == c) return p;
		p++;
	}
	return NULL;
}

__device__ int StrStr(char *src, int src_size, const char *sub, int sub_size)
{
	if (sub_size > src_size) return -1;

	for (int i = 0; i < src_size; i++)
	{
		char *src1 = (char*)&src[i];
		
		if (src_size - i - 1 < sub_size) return -1;

		if (StrCmp(src1, sub_size, sub, sub_size) == 0) return i;
	}

	return -1;
}

__device__ void *MemSet(void *buf, int c, size_t bufsize)
{
	char4 *buf1 = (char4*)buf;

	int len4 = bufsize >> 2;

	int i;
	for (i = 0; i < len4; i++)
	{
		buf1[i].x = c;
		buf1[i].y = c;
		buf1[i].z = c;
		buf1[i].w = c;
	}

	int remainder = bufsize & 3;
	char *rbuf = (char*)&buf1[i];
	for (int i = 0; i < remainder; i++)
		rbuf[i] = c;

	return buf;
}

__device__ int	StrHash(char *str, size_t strLen)
{
	int	hash = strLen;
  
	for (int i=0; i < strLen; i++)
		hash = (hash<<4)^(hash>>28)^str[i];
 
	return hash;
}

__device__ int operator >(char4 str1, char4 str2)
{
	if (str1.x > str2.x) return 0;
	if (str1.y > str2.y) return 0;
	if (str1.z > str2.z) return 0;
	if (str1.w > str2.w) return 0;

	return -1;
}

__device__ int operator <(char4 str1, char4 str2)
{
	if (str1.x < str2.x) return 0;
	if (str1.y < str2.y) return 0;
	if (str1.z < str2.z) return 0;
	if (str1.w < str2.w) return 0;

	return -1;
}

#endif //__BEN_CUDA_STR_CU__