思享工具箱

【CUDA并行编程之八】Cuda实现Kmeans算法

摘要：
本文主要介绍如何使用CUDA并行计算框架编程实现机器学习中的Kmeans算法，Kmeans算法的详细介绍在这里，本文重点在并行实现的过程。free_memory()：释放内存空间~KMEANS()：析构函数并行的代码主要三个函数：find_nearest_cluster(...)compute_delta(...)euclid_dist_2(...)首先看一下函数euclid_dist_2(...)：[cpp]viewplaincopy__host____device__inlinestaticfloateuclid_dist_2{inti;floatans=0;for{ans+=*;}returnans;}这段代码实际上就是并行的计算向量objects[objectId]和clusters[clusterId]之间的距离，即第objectId个数据点到第clusterId个中心点的距离。

本文主要介绍如何使用CUDA并行计算框架编程实现机器学习中的Kmeans算法，Kmeans算法的详细介绍在这里，本文重点在并行实现的过程。

当然还是简单的回顾一下kmeans算法的串行过程：

伪代码：

  1. 创建k个点作为起始质心(经常是随机选择)
  2. 当任意一个点的簇分配结果发生改变时
  3. 对数据集中的每个数据点
  4. 对每个质心
  5. 计算质心与数据点之间的距离
  6. 将数据点分配到距其最近的簇
  7. 对每一个簇，计算簇中所有点的均值并将均值作为质心
我们可以观察到有两个部分可以并行优化：

①line03-04：将每个数据点到多个质心的距离计算进行并行化

②line05：将数据点到某个执行的距离计算进行并行化


KMEANS类：

  1. classKMEANS
  2. {
  3. private:
  4. intnumClusters;
  5. intnumCoords;
  6. intnumObjs;
  7. int*membership;//[numObjs]
  8. char*filename;
  9. float**objects;//[numObjs][numCoords]dataobjects
  10. float**clusters;//[numClusters][unmCoords]clustercenter
  11. floatthreshold;
  12. intloop_iterations;
  14. public:
  15. KMEANS(intk);
  16. voidfile_read(char*fn);
  17. voidfile_write();
  18. voidcuda_kmeans();
  19. inlineintnextPowerOfTwo(intn);
  20. voidfree_memory();
  21. virtual~KMEANS();
  22. };//KMEANS

成员变量：

numClusters：中心点的个数

numCoords：每个数据点的维度

numObjs：数据点的个数

membership：每个数据点所属类别的数组，维度为numObjs

filename：读入的文件名

objects：所有数据点，维度为[numObjs][numCoords]

clusters：中心点数据，维度为[numObjs][numCoords]

threshold：控制循环次数的一个域值

loop_iterations：循环的迭代次数

成员函数：

KMEANS(int k)：含参构造函数。初始化成员变量

file_read(char *fn)：读入文件数据并初始化object以及membership变量

file_write()：将计算结果写回到结果文件中去

cuda_kmeans()：kmeans计算的入口函数

nextPowerOfTwo(int n)：它计算大于等于输入参数n的第一个2的幂次数。

free_memory()：释放内存空间

~KMEANS()：析构函数


并行的代码主要三个函数：

find_nearest_cluster(...)

compute_delta(...)

euclid_dist_2(...)


首先看一下函数euclid_dist_2(...)：

  1. __host____device__inlinestatic
  2. floateuclid_dist_2(intnumCoords,intnumObjs,intnumClusters,float*objects,float*clusters,intobjectId,intclusterId)
  3. {
  4. inti;
  5. floatans=0;
  6. for(i=0;i<numCoords;i++)
  7. {
  8. ans+=(objects[numObjs*i+objectId]-clusters[numClusters*i+clusterId])*
  9. (objects[numObjs*i+objectId]-clusters[numClusters*i+clusterId]);
  10. }
  11. returnans;
  12. }
这段代码实际上就是并行的计算向量objects[objectId]和clusters[clusterId]之间的距离，即第objectId个数据点到第clusterId个中心点的距离。


再看一下函数compute_delta(...)：

  1. /*
  2. *numIntermediates:Theactualnumberofintermediates
  3. *numIntermediates2:Thenextpoweroftwo
  4. */
  5. __global__staticvoidcompute_delta(int*deviceIntermediates,intnumIntermediates,intnumIntermediates2)
  6. {
  7. extern__shared__unsignedintintermediates[];
  9. intermediates[threadIdx.x]=(threadIdx.x<numIntermediates)?deviceIntermediates[threadIdx.x]:0;
  10. __syncthreads();
  12. //numIntermediates2*must*beapoweroftwo!
  13. for(unsignedints=numIntermediates2/2;s>0;s>>=1)
  14. {
  15. if(threadIdx.x<s)
  16. {
  17. intermediates[threadIdx.x]+=intermediates[threadIdx.x+s];
  18. }
  19. __syncthreads();
  20. }
  21. if(threadIdx.x==0)
  22. {
  23. deviceIntermediates[0]=intermediates[0];
  24. }
  25. }
这段代码的意义就是将一个线程块中每个线程的对应的intermediates的数据求和最后放到deviceIntermediates[0]中去然后拷贝回主存块中去。这个问题的更好的解释在这里，实际上就是一个数组求和的问题，应用在这里求得的是有改变的membership中所有数据的和，即改变了簇的点的个数。


最后再看函数finid_nearest_cluster(...)：

  1. /*
  2. *objects:[numCoords][numObjs]
  3. *deviceClusters:[numCoords][numClusters]
  4. *membership:[numObjs]
  5. */
  6. __global__staticvoidfind_nearest_cluster(intnumCoords,intnumObjs,intnumClusters,float*objects,float*deviceClusters,int*membership,int*intermediates)
  7. {
  8. extern__shared__charsharedMemory[];
  9. unsignedchar*membershipChanged=(unsignedchar*)sharedMemory;
  10. float*clusters=deviceClusters;
  12. membershipChanged[threadIdx.x]=0;
  14. intobjectId=blockDim.x*blockIdx.x+threadIdx.x;
  15. if(objectId<numObjs)
  16. {
  17. intindex;
  18. floatdist,min_dist;
  19. /*findtheclusteridthathasmindistancetoobject*/
  20. index=0;
  21. min_dist=euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,0);
  23. for(inti=0;i<numClusters;i++)
  24. {
  25. dist=euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,i);
  26. /*noneedsquareroot*/
  27. if(dist<min_dist)
  28. {
  29. min_dist=dist;
  30. index=i;
  31. }
  32. }
  34. if(membership[objectId]!=index)
  35. {
  36. membershipChanged[threadIdx.x]=1;
  37. }
  38. //assignthemembershiptoobjectobjectId
  39. membership[objectId]=index;
  41. __syncthreads();//formembershipChanged[]
  43. #if1
  44. //blockDim.x*must*beapoweroftwo!
  45. for(unsignedints=blockDim.x/2;s>0;s>>=1)
  46. {
  47. if(threadIdx.x<s)
  48. {
  49. membershipChanged[threadIdx.x]+=membershipChanged[threadIdx.x+s];//calculateallchangedvaluesandsaveresulttomembershipChanged[0]
  50. }
  51. __syncthreads();
  52. }
  53. if(threadIdx.x==0)
  54. {
  55. intermediates[blockIdx.x]=membershipChanged[0];
  56. }
  57. #endif
  58. }
  59. }//find_nearest_cluster
这个函数计算的就是第objectId个数据点到numClusters个中心点的距离，然后根据情况比较更新membership。


这三个函数将所有能够并行的地方都进行了并行，实现了整体算法的并行化~


在此呈上全部代码：

kmeans.h:

  1. #ifndef_H_KMEANS
  2. #define_H_KMEANS
  4. #include<assert.h>
  6. #definemalloc2D(name,xDim,yDim,type)do{
  7. name=(type**)malloc(xDim*sizeof(type*));
  8. assert(name!=NULL);
  9. name[0]=(type*)malloc(xDim*yDim*sizeof(type));
  10. assert(name[0]!=NULL);
  11. for(size_ti=1;i<xDim;i++)
  12. name[i]=name[i-1]+yDim;
  13. }while(0)
  16. doublewtime(void);
  18. #endif

wtime.cu:

  1. #include<sys/time.h>
  2. #include<stdio.h>
  3. #include<stdlib.h>
  5. doublewtime(void)
  6. {
  7. doublenow_time;
  8. structtimevaletstart;
  9. structtimezonetzp;
  11. if(gettimeofday(&etstart,&tzp)==-1)
  12. perror("Error:callinggettimeofday()notsuccessful. ");
  14. now_time=((double)etstart.tv_sec)+/*inseconds*/
  15. ((double)etstart.tv_usec)/1000000.0;/*inmicroseconds*/
  16. returnnow_time;
  17. }


cuda_kmeans.cu：

  1. #include<stdio.h>
  2. #include<stdlib.h>
  3. #include<string.h>
  4. #include<sys/types.h>
  5. #include<sys/stat.h>
  6. #include<unistd.h>
  7. #include<iostream>
  8. #include<cassert>
  10. #include"kmeans.h"
  12. usingnamespacestd;
  14. constintMAX_CHAR_PER_LINE=1024;
  16. classKMEANS
  17. {
  18. private:
  19. intnumClusters;
  20. intnumCoords;
  21. intnumObjs;
  22. int*membership;//[numObjs]
  23. char*filename;
  24. float**objects;//[numObjs][numCoords]dataobjects
  25. float**clusters;//[numClusters][unmCoords]clustercenter
  26. floatthreshold;
  27. intloop_iterations;
  29. public:
  30. KMEANS(intk);
  31. voidfile_read(char*fn);
  32. voidfile_write();
  33. voidcuda_kmeans();
  34. inlineintnextPowerOfTwo(intn);
  35. voidfree_memory();
  36. virtual~KMEANS();
  37. };
  39. KMEANS::~KMEANS()
  40. {
  41. free(membership);
  42. free(clusters[0]);
  43. free(clusters);
  44. free(objects[0]);
  45. free(objects);
  46. }
  48. KMEANS::KMEANS(intk)
  49. {
  50. threshold=0.001;
  51. numObjs=0;
  52. numCoords=0;
  53. numClusters=k;
  54. filename=NULL;
  55. loop_iterations=0;
  56. }
  58. voidKMEANS::file_write()
  59. {
  60. FILE*fptr;
  61. charoutFileName[1024];
  63. //output:thecoordinatesoftheclustercentres
  64. sprintf(outFileName,"%s.cluster_centres",filename);
  65. printf("WritingcoordinatesofK=%dclustercenterstofile"%s" ",numClusters,outFileName);
  66. fptr=fopen(outFileName,"w");
  67. for(inti=0;i<numClusters;i++)
  68. {
  69. fprintf(fptr,"%d",i);
  70. for(intj=0;j<numCoords;j++)
  71. fprintf(fptr,"%f",clusters[i][j]);
  72. fprintf(fptr," ");
  73. }
  74. fclose(fptr);
  76. //output:theclosestclustercentretoeachofthedatapoints
  77. sprintf(outFileName,"%s.membership",filename);
  78. printf("writingmembershipofN=%ddataobjectstofile"%s" ",numObjs,outFileName);
  79. fptr=fopen(outFileName,"w");
  80. for(inti=0;i<numObjs;i++)
  81. {
  82. fprintf(fptr,"%d%d ",i,membership[i]);
  83. }
  84. fclose(fptr);
  85. }
  87. inlineintKMEANS::nextPowerOfTwo(intn)
  88. {
  89. n--;
  90. n=n>>1|n;
  91. n=n>>2|n;
  92. n=n>>4|n;
  93. n=n>>8|n;
  94. n=n>>16|n;
  95. //n=n>>32|n;//for64-bitints
  96. return++n;
  97. }
  99. __host____device__inlinestatic
  100. floateuclid_dist_2(intnumCoords,intnumObjs,intnumClusters,float*objects,float*clusters,intobjectId,intclusterId)
  101. {
  102. inti;
  103. floatans=0;
  104. for(i=0;i<numCoords;i++)
  105. {
  106. ans+=(objects[numObjs*i+objectId]-clusters[numClusters*i+clusterId])*
  107. (objects[numObjs*i+objectId]-clusters[numClusters*i+clusterId]);
  108. }
  109. returnans;
  110. }
  112. /*
  113. *numIntermediates:Theactualnumberofintermediates
  114. *numIntermediates2:Thenextpoweroftwo
  115. */
  116. __global__staticvoidcompute_delta(int*deviceIntermediates,intnumIntermediates,intnumIntermediates2)
  117. {
  118. extern__shared__unsignedintintermediates[];
  120. intermediates[threadIdx.x]=(threadIdx.x<numIntermediates)?deviceIntermediates[threadIdx.x]:0;
  121. __syncthreads();
  123. //numIntermediates2*must*beapoweroftwo!
  124. for(unsignedints=numIntermediates2/2;s>0;s>>=1)
  125. {
  126. if(threadIdx.x<s)
  127. {
  128. intermediates[threadIdx.x]+=intermediates[threadIdx.x+s];
  129. }
  130. __syncthreads();
  131. }
  132. if(threadIdx.x==0)
  133. {
  134. deviceIntermediates[0]=intermediates[0];
  135. }
  136. }
  138. /*
  139. *objects:[numCoords][numObjs]
  140. *deviceClusters:[numCoords][numClusters]
  141. *membership:[numObjs]
  142. */
  143. __global__staticvoidfind_nearest_cluster(intnumCoords,intnumObjs,intnumClusters,float*objects,float*deviceClusters,int*membership,int*intermediates)
  144. {
  145. extern__shared__charsharedMemory[];
  146. unsignedchar*membershipChanged=(unsignedchar*)sharedMemory;
  147. float*clusters=deviceClusters;
  149. membershipChanged[threadIdx.x]=0;
  151. intobjectId=blockDim.x*blockIdx.x+threadIdx.x;
  152. if(objectId<numObjs)
  153. {
  154. intindex;
  155. floatdist,min_dist;
  156. /*findtheclusteridthathasmindistancetoobject*/
  157. index=0;
  158. min_dist=euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,0);
  160. for(inti=0;i<numClusters;i++)
  161. {
  162. dist=euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,i);
  163. /*noneedsquareroot*/
  164. if(dist<min_dist)
  165. {
  166. min_dist=dist;
  167. index=i;
  168. }
  169. }
  171. if(membership[objectId]!=index)
  172. {
  173. membershipChanged[threadIdx.x]=1;
  174. }
  175. //assignthemembershiptoobjectobjectId
  176. membership[objectId]=index;
  178. __syncthreads();//formembershipChanged[]
  180. #if1
  181. //blockDim.x*must*beapoweroftwo!
  182. for(unsignedints=blockDim.x/2;s>0;s>>=1)
  183. {
  184. if(threadIdx.x<s)
  185. {
  186. membershipChanged[threadIdx.x]+=membershipChanged[threadIdx.x+s];//calculateallchangedvaluesandsaveresulttomembershipChanged[0]
  187. }
  188. __syncthreads();
  189. }
  190. if(threadIdx.x==0)
  191. {
  192. intermediates[blockIdx.x]=membershipChanged[0];
  193. }
  194. #endif
  195. }
  196. }//find_nearest_cluster
  198. voidKMEANS::cuda_kmeans()
  199. {
  200. intindex,loop=0;
  201. int*newClusterSize;//[numClusters]:no.objectsassignedineachnewcluster
  202. floatdelta;//%ofobjectschangestheirclusters
  203. float**dimObjects;//[numCoords][numObjs]
  204. float**dimClusters;
  205. float**newClusters;//[numCoords][numClusters]
  207. float*deviceObjects;//[numCoords][numObjs]
  208. float*deviceClusters;//[numCoords][numclusters]
  209. int*deviceMembership;
  210. int*deviceIntermediates;
  212. //Copyobjectsgivenin[numObjs][numCoords]layouttonew[numCoords][numObjs]layout
  213. malloc2D(dimObjects,numCoords,numObjs,float);
  214. for(inti=0;i<numCoords;i++)
  215. {
  216. for(intj=0;j<numObjs;j++)
  217. {
  218. dimObjects[i][j]=objects[j][i];
  219. }
  220. }
  221. //pickfirstnumClusterselementsofobjects[]asinitialclustercenters
  222. malloc2D(dimClusters,numCoords,numClusters,float);
  223. for(inti=0;i<numCoords;i++)
  224. {
  225. for(intj=0;j<numClusters;j++)
  226. {
  227. dimClusters[i][j]=dimObjects[i][j];
  228. }
  229. }
  230. newClusterSize=newint[numClusters];
  231. assert(newClusterSize!=NULL);
  232. malloc2D(newClusters,numCoords,numClusters,float);
  233. memset(newClusters[0],0,numCoords*numClusters*sizeof(float));
  235. //Tosupportreduction,numThreadsPerClusterBlock*must*beapoweroftwo,andit*must*benolargerthanthenumberofbitsthatwillfitintoanunsignedchar,thetypeusedtokeeptrackofmembershipchangesinthekernel.
  236. constunsignedintnumThreadsPerClusterBlock=32;
  237. constunsignedintnumClusterBlocks=(numObjs+numThreadsPerClusterBlock-1)/numThreadsPerClusterBlock;
  238. constunsignedintnumReductionThreads=nextPowerOfTwo(numClusterBlocks);
  240. constunsignedintclusterBlockSharedDataSize=numThreadsPerClusterBlock*sizeof(unsignedchar);
  242. constunsignedintreductionBlockSharedDataSize=numReductionThreads*sizeof(unsignedint);
  244. cudaMalloc(&deviceObjects,numObjs*numCoords*sizeof(float));
  245. cudaMalloc(&deviceClusters,numClusters*numCoords*sizeof(float));
  246. cudaMalloc(&deviceMembership,numObjs*sizeof(int));
  247. cudaMalloc(&deviceIntermediates,numReductionThreads*sizeof(unsignedint));
  249. cudaMemcpy(deviceObjects,dimObjects[0],numObjs*numCoords*sizeof(float),cudaMemcpyHostToDevice);
  250. cudaMemcpy(deviceMembership,membership,numObjs*sizeof(int),cudaMemcpyHostToDevice);
  252. do
  253. {
  254. cudaMemcpy(deviceClusters,dimClusters[0],numClusters*numCoords*sizeof(float),cudaMemcpyHostToDevice);
  256. find_nearest_cluster<<<numClusterBlocks,numThreadsPerClusterBlock,clusterBlockSharedDataSize>>>(numCoords,numObjs,numClusters,deviceObjects,deviceClusters,deviceMembership,deviceIntermediates);
  258. cudaDeviceSynchronize();
  260. compute_delta<<<1,numReductionThreads,reductionBlockSharedDataSize>>>(deviceIntermediates,numClusterBlocks,numReductionThreads);
  262. cudaDeviceSynchronize();
  264. intd;
  265. cudaMemcpy(&d,deviceIntermediates,sizeof(int),cudaMemcpyDeviceToHost);
  266. delta=(float)d;
  268. cudaMemcpy(membership,deviceMembership,numObjs*sizeof(int),cudaMemcpyDeviceToHost);
  270. for(inti=0;i<numObjs;i++)
  271. {
  272. //findthearrayindexofnestest
  273. index=membership[i];
  274. //updatenewclustercenters:sumofobjectslocatedwithin
  275. newClusterSize[index]++;
  276. for(intj=0;j<numCoords;j++)
  277. {
  278. newClusters[j][index]+=objects[i][j];
  279. }
  280. }
  281. //averagethesumandreplaceoldclustercenterswithnewClusters
  282. for(inti=0;i<numClusters;i++)
  283. {
  284. for(intj=0;j<numCoords;j++)
  285. {
  286. if(newClusterSize[i]>0)
  287. dimClusters[j][i]=newClusters[j][i]/newClusterSize[i];
  288. newClusters[j][i]=0.0;//setbackto0
  289. }
  290. newClusterSize[i]=0;//setbackto0
  291. }
  292. delta/=numObjs;
  293. }while(delta>threshold&&loop++<500);
  295. loop_iterations=loop+1;
  297. malloc2D(clusters,numClusters,numCoords,float);
  298. for(inti=0;i<numClusters;i++)
  299. {
  300. for(intj=0;j<numCoords;j++)
  301. {
  302. clusters[i][j]=dimClusters[j][i];
  303. }
  304. }
  306. cudaFree(deviceObjects);
  307. cudaFree(deviceClusters);
  308. cudaFree(deviceMembership);
  309. cudaFree(deviceMembership);
  311. free(dimObjects[0]);
  312. free(dimObjects);
  313. free(dimClusters[0]);
  314. free(dimClusters);
  315. free(newClusters[0]);
  316. free(newClusters);
  317. free(newClusterSize);
  318. }
  320. voidKMEANS::file_read(char*fn)
  321. {
  323. FILE*infile;
  324. char*line=newchar[MAX_CHAR_PER_LINE];
  325. intlineLen=MAX_CHAR_PER_LINE;
  327. filename=fn;
  328. infile=fopen(filename,"r");
  329. assert(infile!=NULL);
  330. /*findthenumberofobjects*/
  331. while(fgets(line,lineLen,infile))
  332. {
  333. numObjs++;
  334. }
  336. /*findthedimensionofeachobject*/
  337. rewind(infile);
  338. while(fgets(line,lineLen,infile)!=NULL)
  339. {
  340. if(strtok(line," ")!=0)
  341. {
  342. while(strtok(NULL," "))
  343. numCoords++;
  344. break;
  345. }
  346. }
  348. /*allocatespaceforobject[][]andreadallobjcet*/
  349. rewind(infile);
  350. objects=newfloat*[numObjs];
  351. for(inti=0;i<numObjs;i++)
  352. {
  353. objects[i]=newfloat[numCoords];
  354. }
  355. inti=0;
  356. /*readallobject*/
  357. while(fgets(line,lineLen,infile)!=NULL)
  358. {
  359. if(strtok(line," ")==NULL)continue;
  360. for(intj=0;j<numCoords;j++)
  361. {
  362. objects[i][j]=atof(strtok(NULL,", "));
  363. }
  364. i++;
  365. }
  367. /*membership:theclusteridforeachdataobject*/
  368. membership=newint[numObjs];
  369. assert(membership!=NULL);
  370. for(inti=0;i<numObjs;i++)
  371. membership[i]=-1;
  373. }
  375. intmain(intargc,char*argv[])
  376. {
  377. KMEANSkmeans(atoi(argv[1]));
  378. kmeans.file_read(argv[2]);
  379. kmeans.cuda_kmeans();
  380. kmeans.file_write();
  381. return0;
  382. }


makefile：

  1. target:
  2. nvcccuda_kmeans.cu
  3. ./a.out4./Image_data/color100.txt

所有代码和文件数据在这里：http://yunpan.cn/cKBZMPAJ8tcAs（提取码：9476）


运行代码：

【CUDA并行编程之八】Cuda实现Kmeans算法第19张


kmeans的cuda实现代码相对复杂，在阅读的过程中可能会有困难，有问题请留言~


Author：忆之独秀

Email：leaguenew@qq.com

注明出处：http://blog.csdn.net/lavorange/article/details/41942323


标签：#kmeans#char函数#sizeof#char#cuda | 浏览：463 | 发布日期：

免责声明：文章转载自《【CUDA并行编程之八】Cuda实现Kmeans算法》仅用于学习参考。如对内容有疑问，请及时联系本站处理。

上篇[UWP小白日记2]SQLite数据库DOMEBootstrap实现弹出框和提示框效果代码下篇

宿迁高防，2C2G15M，22元/月；香港BGP，2C5G5M，25元/月 雨云优惠码：MjYwNzM=

相关文章

ws2s函数

std::string ws2s(const std::wstring&str) { char*pElementText; intiTextLen; //wide char to multi char iTextLen = WideCharToMultiByte(CP_ACP, 0, str.c_str(), -1,...

C#中char[]与string之间的转换；byte[]与string之间的转化

（1）C#中char[]与string互相转换的写法：string 转换成 Char[]string ss="abcdefg";char[] cc=ss.ToCharArray();Char[] 转换成stringstring s=new string(cc);------------------------------------------------...

C++获取MAC与IP

#include <Nb30.h> #pragma comment(lib,"ws2_32.lib") #pragma comment(lib,"netapi32.lib") std::stringGetMac() { NCB ncb; typedef struct_ASTAT_ { ADA...

【转载】C/C++中的char，wchar，TCHAR

点击这里查看原文章 总体简介：由于字符编码的不同，在C++中有三种对于字符类型：char, wchar_t , TCHAR。其实TCHAR不能算作一种类型，他紧紧是一个宏。我们都知道，宏在预编译的时候会被替换成相应的内容。TCHAR 在使用多字节编码时被定义成char，在Unicode编码时定义成wchar_t。 1.VC++中的char,wchar_t...

DAG优化

DAG优化 Time Limit: 1000 ms Memory Limit: 65536 KiB Problem Description 大家都学过了代码优化，其中有一个DAG优化，这次我们就练习这个操作。 Input 输入第一行为一个整数ｎ(n < 100)，表示该组输入的表达式的个数...

CAS原子操作实现无锁及性能分析

  CAS原子操作实现无锁及性能分析 Author：Echo Chen（陈斌） Email：chenb19870707@gmail.com Blog：Blog.csdn.net/chen19870707 Date：Nov 13th, 2014 近期在研究nginx的自旋锁的时候，又见到了GCC CAS原子操作，于是决定动手分析下CAS实现的无锁究竟性能...

最新文章

随机推荐

思享工具箱导航