6月28日更新--“如何并行化我的应用”方案有奖征集活动--获奖名单公布

图腾部落

想参赛的朋友请抓紧时间，本活动参赛作品提交截止日期为6月15日。

yu_single

请教不敢当啊，十分欢迎交流切磋哈；）

bewarm

围观一下...不说话

cyrosly

我就直接发代码参赛吧，多少钱随便给

图腾部落

原帖由 cyrosly 于 2010-6-7 08:46 发表
我就直接发代码参赛吧，多少钱随便给

别纯代码，配点说明，至少项目背景大概说下。

lloves

K近邻算法基于CUDA的实现

前面看到有个人写了个k-均值算法，我今天就写个ｋ近邻算法看看！

KNN 算法的基本思想是:将文本内容形式化为特征空间中的向量,每个文本就是一个向量，向量的每维数据代表不同的内容。对于一个测试文本, 计算它与已分类样本集中每个文本的相似度, 找出K 个最相似的文本, 根据这k个文本中属于那个类的文本最多，以此判断测试文本所属的类别。具体算法步骤如下:
( 1) 将待分类文本特征化为向量a
( 2) 计算该测试文本与训练集中每个文本的文本（特征化为向量ｂ）相似度, 计算公式为:　cos = ab/(|a||b|)
( 3) 按照文本相似度, 在训练文本集中选出与测试文本最相似的k 个文本。
( 4) 在测试文本的k 个近邻中, 依次计算每类的权重。本文权重定义为文本数
( 5) 比较类的权重, 将文本分到权重最大的那个类别中。

CUDA实现

1. 
   
2. typedef struct text{
   
3.         unsigned int len;//text number
   
4.         unsigned int dim;//text’s dimension
   
5. //every text have a row, column first
   
6.         float *textData;
   
7.         unsigned int pitch;//use to access textData
   
8. }text;
   

复制代码

text结构中，len　表示所有文本的数量；dim　表示每个文本向量的维数；textData　表示所有的文本数据，每列代表一个文本的特征信息，pitch　是用cudaMallocPitch分配数据时返回的值除以４，也就是textData每行的长度。
下面的代码用于计算两个文本的相似度

1. 
   
2. //first ,second : text index
   
3. __host__ __device__ float homo(text t, int first, int second){
   
4.         float mul = 0.0f;
   
5.         float o2 = 0.0f, tw2 = 0.0f, temp1, temp2;
   
6.        
   
7.         for(int I = 0; I < t.dim; i++){
   
8.                 temp1 = t.textData[first + t.pitch*i];
   
9.                 temp2 = t.textData[second + t.pitch*i];
   
10.                 mul += temp1*temp2;
    
11.                 o2 +=temp1*temp1;
    
12.                 tw2 += temp2*temp2;
    
13. }
    
14. 
    
15.         return mul/sqrtf(o2*tw2);
    
16. }
    

复制代码

内核函数大致如下：其基本思想相当简单，就是每个线程处理自己的那一亩三分地。

1. 
   
2. __global__ void knnKernel(Text t,int k, int * d_belongTo){
   
3.         Int id = blockDim.x*blockIdx.x + threadIdx.x;
   
4. 
   
5.         float knn[K];
   
6.         int index[K];
   
7. 
   
8.         for(int I =0; I < k; i++){
   
9.                 index[i] = 0;
   
10.                 knn[i]= 0.0f;
    
11.         ｝
    
12. 
    
13.         If( id >= t.len)
    
14.                 return;
    
15. 
    
16.         float temp =home(…);
    
17. //遍历knn，将ｔｅｍｐ插入其中，丢弃最小的那个并记录其所属的类
    
18. 
    
19. …………………
    
20. 
    
21. }
    

复制代码

由于流多处理器中的寄存器数量有限，所以Ｋ值不能太大，具体大小依据硬件每个 SM寄存器的数目以及每个线程块的线程数决定，由于 CUDA 目前还不支持动态数据，因此只能使用宏,这样保证了在实际使用的k值小于预定义的K时，依旧会被分配到寄存器中，但是一旦k大于K时，就要重新处理了。

[ 本帖最后由 lloves 于 2010-6-7 19:21 编辑 ]

图腾部落

原帖由 lloves 于 2010-6-7 19:18 发表
K近邻算法基于CUDA的实现

前面看到有个人写了个k-均值算法，我今天就写个ｋ近邻算法看看！

KNN 算法的基本思想是:将文本内容形式化为特征空间中的向量,每个文本就是一个向量，向量的每维数据代表不同的内容。对于一个测试文本, 计算它与已分类样本集中每个文本的相似度, 找出K 个最相似的文本, 根据这k个文本中属于那个类的文本最多，以此判断测试文本所属的类别。具体算法步骤如下:
( 1) 将待分类文本特征化为向量a
( 2) 计算该测试文本与训练集中每个文本的文本（特征化为向量ｂ）相似度, 计算公式为:　cos = ab/(|a||b|)
( 3) 按照文本相似度, 在训练文本集中选出与测试文本最相似的k 个文本。
( 4) 在测试文本的k 个近邻中, 依次计算每类的权重。本文权重定义为文本数
( 5) 比较类的权重, 将文本分到权重最大的那个类别中。

CUDA实现

typedef struct text{
        unsigned int len;//text number
        unsigned int dim;//text’s dimension
//every text have a row, column first
        float *textData;
        unsigned int pitch;//use to access textData
}text;


text结构中，len　表示所有文本的数量；dim　表示每个文本向量的维数；textData　表示所有的文本数据，每列代表一个文本的特征信息，pitch　是用cudaMallocPitch分配数据时返回的值除以４，也就是textData每行的长度。
下面的代码用于计算两个文本的相似度

//first ,second : text index
__host__ __device__ float homo(text t, int first, int second){
        float mul = 0.0f;
        float o2 = 0.0f, tw2 = 0.0f, temp1, temp2;
       
        for(int I = 0; I < t.dim; i++){
                temp1 = t.textData[first + t.pitch*i];
                temp2 = t.textData[second + t.pitch*i];
                mul += temp1*temp2;
                o2 +=temp1*temp1;
                tw2 += temp2*temp2;
}

        return mul/sqrtf(o2*tw2);
}


内核函数大致如下：其基本思想相当简单，就是每个线程处理自己的那一亩三分地。

__global__ void knnKernel(Text t,int k, int * d_belongTo){
        Int id = blockDim.x*blockIdx.x + threadIdx.x;

        float knn[K];
        int index[K];

        for(int I =0; I < k; i++){
                index = 0;
                knn= 0.0f;
        ｝

        If( id >= t.len)
                return;

        float temp =home(…);
//遍历knn，将ｔｅｍｐ插入其中，丢弃最小的那个并记录其所属的类

…………………

}

由于流多处理器中的寄存器数量有限，所以Ｋ值不能太大，具体大小依据硬件每个 SM寄存器的数目以及每个线程块的线程数决定，由于 CUDA 目前还不支持动态数据，因此只能使用宏,这样保证了在实际使用的k值小于预定义的K时，依旧会被分配到寄存器中，但是一旦k大于K时，就要重新处理了。

cyrosly

Be cleared

[ 本帖最后由 cyrosly 于 2010-6-30 10:22 编辑 ]

yyfn风辰

有时间研究一下你的代码，呵呵！

lloves