cuda string library

 in this article, i will write my note on cuda string library which is written by a professor in the hkust.

 

first, let us look at strcpy implemented in cuda.

__device__ char* strcpy(char * dst, char * src,int size) { char4 * dst1 = (char4*)dst; char4 * src1 = (char4*)src; //because we use char4 , so it is necessary to divide the number of characters by 4. int numChar4 = size>>2; int i = 0; for(;i < numChar4,++i) { dst[i] = src[i]; } //maybe the size can not be divided by 4 exactly,so we need copy the remained characters. for(i = (numChar4<<2);i<size;++i) { dst1[i] = src1[i]; } }

 in cuda,it is recommended to use build-in vector like char4 and int4, because gpu can fetch the entire vector in one clock cycle from the memory.

 __device__ int strcmp(char * str1, char * str2, int size1,int size2) { //first compare the minsize characters. int minsize = size1 < size2? size1:size2; char4* char4_str1 = (char4*)str1; char4* char4_str2 = (char4*)str2; //same with strcpy,each time we read 4 characters,so charnumsize should be divided by 4 int charnumsize = minsize >> 2; int i = 0 ; char4 s1, s2;//build-in vector which has 4 elements x,y,z,w. for(; i < charnumsize;++i) { s1 = char4_str1[i]; s2 = char4_str2[i]; //comparasion begins with the first character and then the second,third and fouth one. if(s1.x > s2.x)return 1; else if(s1.x<s2.x)return -1; if(s1.y > s2.y) return 1; else if(s1.y < s2.y)return -1; if(s1.z > s2.z)return 1; else if(s1.z < s2.z)return -1; if(s1.w > s2.w)return 1; else if(s1.w < s2.w) return -1; } //compare the remained characters. for(i = (charnumsize << 2); i < minsize;++i) { if(str1[i] > str2[i] )return 1; else if(str1[i] < str2[i])return -1; } //if the minsize characters are equal then we should compare their length. if(size1 > size2)return 1; else if(size1 < size2 )return -1; else return 0; }

strcmp is very like strcpy. it uses build-in vector char4  to speed up memory fetch.

it is a good idea to first compare the first n characters , where n is the length of the shorter string because we don't need to always worry about the end mark '/0' of the string.

if the first n characters of two strings are equal then we just need to compare their length to get the result.

 

additionally, we also can use this method in the memset.look at the code as follows.

__device__ void * memset(char * str, int value , int size) { char4 * str1 = (char*)str; //the same with strcpy int charnum = size >> 2; int i = 0 ; for(;i < charnum ; ++i) { str1[i].x = value; str1[i].y = value; str1[i].z = value; str1[i].w = value; } //size & 3 = 0 , 1, 2 ,or 3 int remainer = size & 3; //rbuf points to the start address of the remained characters. char * rbuf = (char*)(&str[i]); for(i = 0 ; i< remainer; ++i) { rbuf[i] = str[i]; } }