zoukankan      html  css  js  c++  java

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


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

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

    伪代码：

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

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

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


    KMEANS类：

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. class KMEANS  
    2. {  
    3. private:  
    4.     int numClusters;  
    5.     int numCoords;  
    6.     int numObjs;  
    7.     int *membership;//[numObjs]  
    8.     char *filename;   
    9.     float **objects;//[numObjs][numCoords] data objects  
    10.     float **clusters;//[numClusters][unmCoords] cluster center  
    11.     float threshold;  
    12.     int loop_iterations;  
    13.   
    14. public:  
    15.     KMEANS(int k);  
    16.     void file_read(char *fn);  
    17.     void file_write();  
    18.     void cuda_kmeans();  
    19.     inline int nextPowerOfTwo(int n);  
    20.     void free_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(...)：

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. __host__ __device__ inline static   
    2. float euclid_dist_2(int numCoords,int numObjs,int numClusters,float *objects,float *clusters,int objectId,int clusterId)  
    3. {  
    4.     int i;  
    5.     float ans = 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.     return ans;  
    12. }  
    这段代码实际上就是并行的计算向量objects[objectId]和clusters[clusterId]之间的距离，即第objectId个数据点到第clusterId个中心点的距离。


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

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. /* 
    2. * numIntermediates:The actual number of intermediates 
    3. * numIntermediates2:The next power of two 
    4. */  
    5. __global__ static void compute_delta(int *deviceIntermediates,int numIntermediates, int numIntermediates2)  
    6. {  
    7.     extern __shared__ unsigned int intermediates[];  
    8.   
    9.     intermediates[threadIdx.x] = (threadIdx.x < numIntermediates) ? deviceIntermediates[threadIdx.x] : 0 ;  
    10.     __syncthreads();  
    11.   
    12.     //numIntermediates2 *must* be a power of two!  
    13.     for(unsigned int s = 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(...)：

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. /* 
    2. * objects:[numCoords][numObjs] 
    3. * deviceClusters:[numCoords][numClusters] 
    4. * membership:[numObjs] 
    5. */  
    6. __global__ static void find_nearest_cluster(int numCoords,int numObjs,int numClusters,float *objects, float *deviceClusters,int *membership ,int *intermediates)  
    7. {  
    8.     extern __shared__ char sharedMemory[];  
    9.     unsigned char *membershipChanged = (unsigned char *)sharedMemory;  
    10.     float *clusters = deviceClusters;  
    11.   
    12.     membershipChanged[threadIdx.x] = 0;  
    13.   
    14.     int objectId = blockDim.x * blockIdx.x + threadIdx.x;  
    15.     if( objectId < numObjs )  
    16.     {  
    17.         int index;  
    18.         float dist,min_dist;  
    19.         /*find the cluster id that has min distance to object*/  
    20.         index = 0;  
    21.         min_dist = euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,0);  
    22.           
    23.         for(int i=0;i<numClusters;i++)  
    24.         {  
    25.             dist = euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,i) ;  
    26.             /* no need square root */  
    27.             if( dist < min_dist )  
    28.             {  
    29.                 min_dist = dist;  
    30.                 index = i;  
    31.             }  
    32.         }  
    33.   
    34.         if( membership[objectId]!=index )  
    35.         {  
    36.             membershipChanged[threadIdx.x] = 1;   
    37.         }  
    38.         //assign the membership to object objectId  
    39.         membership[objectId] = index;  
    40.   
    41.         __syncthreads(); //for membershipChanged[]  
    42.   
    43. #if 1  
    44.         //blockDim.x *must* be a power of two!  
    45.         for(unsigned int s = blockDim.x / 2; s > 0 ;s>>=1)  
    46.         {  
    47.             if(threadIdx.x < s)    
    48.             {  
    49.                 membershipChanged[threadIdx.x] += membershipChanged[threadIdx.x + s];//calculate all changed values and save result to membershipChanged[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:

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. #ifndef _H_KMEANS  
    2. #define _H_KMEANS  
    3.   
    4. #include <assert.h>  
    5.   
    6. #define malloc2D(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_t i = 1; i < xDim; i++)                         
    12.         name[i] = name[i-1] + yDim;                           
    13. } while (0)  
    14.   
    15.   
    16. double  wtime(void);  
    17.   
    18. #endif  

    wtime.cu:

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. #include <sys/time.h>  
    2. #include <stdio.h>  
    3. #include <stdlib.h>  
    4.   
    5. double wtime(void)   
    6. {  
    7.     double          now_time;  
    8.     struct timeval  etstart;  
    9.     struct timezone tzp;  
    10.   
    11.     if (gettimeofday(&etstart, &tzp) == -1)  
    12.         perror("Error: calling gettimeofday() not successful. ");  
    13.   
    14.     now_time = ((double)etstart.tv_sec) +              /* in seconds */  
    15.                ((double)etstart.tv_usec) / 1000000.0;  /* in microseconds */  
    16.     return now_time;  
    17. }  


    cuda_kmeans.cu：

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    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>  
    9.   
    10. #include "kmeans.h"  
    11.   
    12. using namespace std;  
    13.   
    14. const int MAX_CHAR_PER_LINE = 1024;  
    15.   
    16. class KMEANS  
    17. {  
    18. private:  
    19.     int numClusters;  
    20.     int numCoords;  
    21.     int numObjs;  
    22.     int *membership;//[numObjs]  
    23.     char *filename;   
    24.     float **objects;//[numObjs][numCoords] data objects  
    25.     float **clusters;//[numClusters][unmCoords] cluster center  
    26.     float threshold;  
    27.     int loop_iterations;  
    28.   
    29. public:  
    30.     KMEANS(int k);  
    31.     void file_read(char *fn);  
    32.     void file_write();  
    33.     void cuda_kmeans();  
    34.     inline int nextPowerOfTwo(int n);  
    35.     void free_memory();  
    36.     virtual ~KMEANS();  
    37. };  
    38.   
    39. KMEANS::~KMEANS()  
    40. {  
    41.     free(membership);  
    42.     free(clusters[0]);  
    43.     free(clusters);  
    44.     free(objects[0]);  
    45.     free(objects);  
    46. }  
    47.   
    48. KMEANS::KMEANS(int k)  
    49. {  
    50.     threshold = 0.001;  
    51.     numObjs = 0;  
    52.     numCoords = 0;  
    53.     numClusters = k;  
    54.     filename = NULL;  
    55.     loop_iterations = 0;  
    56. }  
    57.   
    58. void KMEANS::file_write()  
    59. {  
    60.     FILE *fptr;  
    61.     char outFileName[1024];  
    62.   
    63.     //output:the coordinates of the cluster centres  
    64.     sprintf(outFileName,"%s.cluster_centres",filename);  
    65.     printf("Writingcoordinates of K=%d cluster centers to file "%s" ",numClusters,outFileName);  
    66.     fptr = fopen(outFileName,"w");  
    67.     for(int i=0;i<numClusters;i++)  
    68.     {  
    69.         fprintf(fptr,"%d ",i)   ;  
    70.         for(int j=0;j<numCoords;j++)  
    71.             fprintf(fptr,"%f ",clusters[i][j]);  
    72.         fprintf(fptr," ");  
    73.     }  
    74.     fclose(fptr);  
    75.   
    76.     //output:the closest cluster centre to each of the data points  
    77.     sprintf(outFileName,"%s.membership",filename);  
    78.     printf("writing membership of N=%d data objects to file "%s"  ",numObjs,outFileName);  
    79.     fptr = fopen(outFileName,"w");  
    80.     for(int i=0;i<numObjs;i++)  
    81.     {  
    82.         fprintf(fptr,"%d %d ",i,membership[i]) ;  
    83.     }  
    84.     fclose(fptr);  
    85. }  
    86.   
    87. inline int KMEANS::nextPowerOfTwo(int n)  
    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; // for 64-bit ints  
    96.     return ++n;  
    97. }  
    98.   
    99. __host__ __device__ inline static   
    100. float euclid_dist_2(int numCoords,int numObjs,int numClusters,float *objects,float *clusters,int objectId,int clusterId)  
    101. {  
    102.     int i;  
    103.     float ans = 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.     return ans;  
    110. }  
    111.   
    112. /* 
    113. * numIntermediates:The actual number of intermediates 
    114. * numIntermediates2:The next power of two 
    115. */  
    116. __global__ static void compute_delta(int *deviceIntermediates,int numIntermediates, int numIntermediates2)  
    117. {  
    118.     extern __shared__ unsigned int intermediates[];  
    119.   
    120.     intermediates[threadIdx.x] = (threadIdx.x < numIntermediates) ? deviceIntermediates[threadIdx.x] : 0 ;  
    121.     __syncthreads();  
    122.   
    123.     //numIntermediates2 *must* be a power of two!  
    124.     for(unsigned int s = 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. }  
    137.   
    138. /* 
    139. * objects:[numCoords][numObjs] 
    140. * deviceClusters:[numCoords][numClusters] 
    141. * membership:[numObjs] 
    142. */  
    143. __global__ static void find_nearest_cluster(int numCoords,int numObjs,int numClusters,float *objects, float *deviceClusters,int *membership ,int *intermediates)  
    144. {  
    145.     extern __shared__ char sharedMemory[];  
    146.     unsigned char *membershipChanged = (unsigned char *)sharedMemory;  
    147.     float *clusters = deviceClusters;  
    148.   
    149.     membershipChanged[threadIdx.x] = 0;  
    150.   
    151.     int objectId = blockDim.x * blockIdx.x + threadIdx.x;  
    152.     if( objectId < numObjs )  
    153.     {  
    154.         int index;  
    155.         float dist,min_dist;  
    156.         /*find the cluster id that has min distance to object*/  
    157.         index = 0;  
    158.         min_dist = euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,0);  
    159.           
    160.         for(int i=0;i<numClusters;i++)  
    161.         {  
    162.             dist = euclid_dist_2(numCoords,numObjs,numClusters,objects,clusters,objectId,i) ;  
    163.             /* no need square root */  
    164.             if( dist < min_dist )  
    165.             {  
    166.                 min_dist = dist;  
    167.                 index = i;  
    168.             }  
    169.         }  
    170.   
    171.         if( membership[objectId]!=index )  
    172.         {  
    173.             membershipChanged[threadIdx.x] = 1;   
    174.         }  
    175.         //assign the membership to object objectId  
    176.         membership[objectId] = index;  
    177.   
    178.         __syncthreads(); //for membershipChanged[]  
    179.   
    180. #if 1  
    181.         //blockDim.x *must* be a power of two!  
    182.         for(unsigned int s = blockDim.x / 2; s > 0 ;s>>=1)  
    183.         {  
    184.             if(threadIdx.x < s)    
    185.             {  
    186.                 membershipChanged[threadIdx.x] += membershipChanged[threadIdx.x + s];//calculate all changed values and save result to membershipChanged[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  
    197.   
    198. void KMEANS::cuda_kmeans()  
    199. {  
    200.     int index,loop = 0;  
    201.     int *newClusterSize;//[numClusters]:no.objects assigned in each new cluster  
    202.     float delta; //% of objects changes their clusters  
    203.     float **dimObjects;//[numCoords][numObjs]  
    204.     float **dimClusters;  
    205.     float **newClusters;//[numCoords][numClusters]  
    206.   
    207.     float *deviceObjects; //[numCoords][numObjs]  
    208.     float *deviceClusters; //[numCoords][numclusters]  
    209.     int *deviceMembership;  
    210.     int *deviceIntermediates;  
    211.   
    212.     //Copy objects given in [numObjs][numCoords] layout to new [numCoords][numObjs] layout  
    213.     malloc2D(dimObjects,numCoords,numObjs,float);  
    214.     for(int i=0;i<numCoords;i++)  
    215.     {  
    216.         for(int j=0;j<numObjs;j++)  
    217.         {  
    218.             dimObjects[i][j] = objects[j][i];     
    219.         }  
    220.     }  
    221.     //pick first numClusters elements of objects[] as initial cluster centers  
    222.     malloc2D(dimClusters, numCoords, numClusters,float);  
    223.     for(int i=0;i<numCoords;i++)  
    224.     {  
    225.         for(int j=0;j<numClusters;j++)  
    226.         {  
    227.             dimClusters[i][j] = dimObjects[i][j];  
    228.         }  
    229.     }  
    230.     newClusterSize = new int[numClusters];  
    231.     assert(newClusterSize!=NULL);  
    232.     malloc2D(newClusters,numCoords,numClusters,float);  
    233.     memset(newClusters[0],0,numCoords * numClusters * sizeof(float) );  
    234.       
    235.     //To support reduction,numThreadsPerClusterBlock *must* be a power of two, and it *must* be no larger than the number of bits that will fit into an unsigned char ,the type used to keep track of membership changes in the kernel.  
    236.     const unsigned int numThreadsPerClusterBlock = 32;  
    237.     const unsigned int numClusterBlocks = (numObjs + numThreadsPerClusterBlock -1)/numThreadsPerClusterBlock;  
    238.     const unsigned int numReductionThreads = nextPowerOfTwo(numClusterBlocks);  
    239.   
    240.     const unsigned int clusterBlockSharedDataSize = numThreadsPerClusterBlock * sizeof(unsigned char);  
    241.   
    242.     const unsigned int reductionBlockSharedDataSize = numReductionThreads * sizeof(unsigned int);  
    243.   
    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(unsigned int));  
    248.   
    249.     cudaMemcpy(deviceObjects,dimObjects[0],numObjs*numCoords*sizeof(float),cudaMemcpyHostToDevice);  
    250.     cudaMemcpy(deviceMembership,membership,numObjs*sizeof(int),cudaMemcpyHostToDevice);  
    251.   
    252.     do  
    253.     {  
    254.         cudaMemcpy(deviceClusters,dimClusters[0],numClusters*numCoords*sizeof(float),cudaMemcpyHostToDevice);  
    255.   
    256.         find_nearest_cluster<<<numClusterBlocks,numThreadsPerClusterBlock,clusterBlockSharedDataSize>>>(numCoords,numObjs,numClusters,deviceObjects,deviceClusters,deviceMembership,deviceIntermediates);  
    257.   
    258.         cudaDeviceSynchronize();  
    259.   
    260.         compute_delta<<<1,numReductionThreads,reductionBlockSharedDataSize>>>(deviceIntermediates,numClusterBlocks,numReductionThreads);  
    261.   
    262.         cudaDeviceSynchronize();  
    263.           
    264.         int d;  
    265.         cudaMemcpy(&d,deviceIntermediates,sizeof(int),cudaMemcpyDeviceToHost);  
    266.         delta = (float)d;  
    267.   
    268.         cudaMemcpy(membership,deviceMembership,numObjs*sizeof(int),cudaMemcpyDeviceToHost);  
    269.           
    270.         for(int i=0;i<numObjs;i++)  
    271.         {  
    272.             //find the array index of nestest   
    273.             index = membership[i];  
    274.             //update new cluster centers:sum of objects located within  
    275.             newClusterSize[index]++;  
    276.             for(int j=0;j<numCoords;j++)  
    277.             {  
    278.                 newClusters[j][index] += objects[i][j];  
    279.             }  
    280.         }  
    281.         //average the sum and replace old cluster centers with newClusters   
    282.         for(int i=0;i<numClusters;i++)  
    283.         {  
    284.             for(int j=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;//set back to 0  
    289.             }  
    290.             newClusterSize[i] = 0 ; //set back to 0  
    291.         }  
    292.         delta /= numObjs;  
    293.     }while( delta > threshold && loop++ < 500 );  
    294.   
    295.     loop_iterations = loop + 1;  
    296.       
    297.     malloc2D(clusters,numClusters,numCoords,float);  
    298.     for(int i=0;i<numClusters;i++)  
    299.     {  
    300.         for(int j=0;j<numCoords;j++)  
    301.         {  
    302.             clusters[i][j] = dimClusters[j][i];  
    303.         }  
    304.     }  
    305.   
    306.     cudaFree(deviceObjects) ;  
    307.     cudaFree(deviceClusters);  
    308.     cudaFree(deviceMembership);  
    309.     cudaFree(deviceMembership);  
    310.   
    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. }  
    319.   
    320. void KMEANS::file_read(char *fn)  
    321. {  
    322.   
    323.     FILE *infile;  
    324.     char *line = new char[MAX_CHAR_PER_LINE];  
    325.     int lineLen = MAX_CHAR_PER_LINE;  
    326.   
    327.     filename = fn;  
    328.     infile = fopen(filename,"r");  
    329.     assert(infile!=NULL);  
    330.     /*find the number of objects*/    
    331.     while( fgets(line,lineLen,infile) )  
    332.     {  
    333.         numObjs++;    
    334.     }  
    335.   
    336.     /*find the dimension of each object*/  
    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.     }  
    347.   
    348.     /*allocate space for object[][] and read all objcet*/  
    349.     rewind(infile);  
    350.     objects = new float*[numObjs];  
    351.     for(int i=0;i<numObjs;i++)  
    352.     {  
    353.         objects[i] = new float[numCoords];  
    354.     }  
    355.     int i=0;  
    356.     /*read all object*/  
    357.     while( fgets(line,lineLen,infile)!=NULL )  
    358.     {  
    359.         if( strtok(line,"  ") ==NULL ) continue;  
    360.         for(int j=0;j<numCoords;j++)  
    361.         {  
    362.             objects[i][j] = atof( strtok(NULL," , ") )   ;  
    363.         }  
    364.         i++;  
    365.     }  
    366.       
    367.     /* membership: the cluster id for each data object */  
    368.     membership = new int[numObjs];  
    369.     assert(membership!=NULL);  
    370.     for(int i=0;i<numObjs;i++)  
    371.         membership[i] = -1;  
    372.       
    373. }  
    374.   
    375. int main(int argc,char *argv[])  
    376. {  
    377.     KMEANS kmeans(atoi(argv[1]));  
    378.     kmeans.file_read(argv[2]);  
    379.     kmeans.cuda_kmeans();  
    380.     kmeans.file_write();  
    381.     return 0;  
    382. }  


    makefile：

    [cpp] view plaincopy在CODE上查看代码片派生到我的代码片
    1. target:  
    2.     nvcc cuda_kmeans.cu  
    3.     ./a.out  4 ./Image_data/color100.txt  

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


    运行代码：



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


    Author：忆之独秀

    Email：leaguenew@qq.com

    注明出处：http://blog.csdn.net/lavorange/article/details/41942323
  • 相关阅读:
    CI/CD for Power Platform
    SpringMVC异常处理
    SpringMVC框架中的拦截器
    spring实现文件上传
    idea常用的快捷键
    解决maven项目创建过慢的问题
    springmvc—入门程序
    Spring中的 JdbcTemplate
    基于XML的AOP 配置
    基于注解的 AOP 配置
  • 原文地址：https://www.cnblogs.com/walccott/p/4956932.html
Copyright © 2011-2022 走看看