24小时热门版块排行榜    

查看: 471  |  回复: 1

hsfy919

银虫 (初入文坛)

[求助] CUDA并行问题,全部金币悬赏

最近在研究CUDA做并行运算,刚开始接触,不是很明白共享内存(shared memory)的机理
问题是:我要利用CUDA函数做N1个数组相加,每个数组长度为N2,由于数组很多(大概上百万个)不能通过CPU直接传入GPU中,需要在GPU中通过计算每个线程得到一个数组,完成计算后所有线程的数组相加,得到长度为N2的一个数组,输出出来。可由于线程之间彼此不能通信,想用shared memory来解决求和,但结果不对(我编写的程序如下)。

这个程序正确的结果应该是C=[150,300,450,600,750],但得到的结果却是C=[750,750,750,750,750]
希望大家看看,如果能解决问题,全部金币奉上

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <fstream>
using namespace std;

#define NumberOfBlock 128
#define ThreadPerBlock 256

__global__ void addKernel(int *c, int *p, int *a, int *b, int N)
{
    int id = threadIdx.x + blockIdx.x*blockDim.x, cacheindex=threadIdx.x, i, j, ii;
        double temp, temp2;
        __shared__ double cache[ThreadPerBlock];
        for(j=0;j<N;j++)  // 问题在于当j循环到1时,得到的值覆盖了数组C中j=0时的值,当j=2时覆盖了j=0和1的值,后面也是
        {
                i=id;
            temp=0;
                while(i<N)
                {
                        temp+=a*b[j];
                        i+=gridDim.x*blockDim.x;
                }
//----对每个j规约求和--------------------------------------------------------------------
                cache[cacheindex]=temp;
                __syncthreads();
                ii=blockDim.x/2;
                while(ii!=0)
                {
                        if(cacheindex<ii)
                                cache[cacheindex]+=cache[cacheindex+ii];
                                __syncthreads();
                                ii/=2;
                }
                if(cacheindex==0)
                        p[blockIdx.x]=cache[0];
                __syncthreads();
                temp2=0;
                for(ii=0;ii<NumberOfBlock;ii++)
                        temp2+=p[ii];
                __syncthreads();
                c[j]=temp2;
//--------------------------------------------------------------------------------
        }

}

int main()
{
    const int size = 5;  //以5个元素举例,实际问题元素很多,过百万个
    int a[size] = { 1, 2, 3, 4, 5 };
    int b[size] = { 10, 20, 30, 40, 50 };
        int c[size];
        for(int i=0;i<size;i++)
                c = 0;
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
        int *dev_p = 0;

        int *p=new int [NumberOfBlock];
        for(int i=0;i<NumberOfBlock;i++)
                p=0;
   
    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_p, NumberOfBlock * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));
    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_c, c, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_p, p, NumberOfBlock * sizeof(int), cudaMemcpyHostToDevice);

    addKernel<<<NumberOfBlock, ThreadPerBlock>>>(dev_c, dev_p, dev_a, dev_b, size);  
    cudaDeviceSynchronize();

    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    for(int i=0;i<size;i++)
        cout<<c<<endl;
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
        getchar();
}
回复此楼
希望优秀成为一种习惯
已阅   回复此楼   关注TA 给TA发消息 送TA红花 TA的回帖

水博士

木虫 (正式写手)

尝试一下用单线程读写,是否正确
2楼2015-03-09 10:55:38
已阅   回复此楼   关注TA 给TA发消息 送TA红花 TA的回帖
相关版块跳转 我要订阅楼主 hsfy919 的主题更新
最具人气热帖推荐 [查看全部] 作者 回/看 最后发表
[基金申请] 售SCI一区文章,我:8O5.5.1.O5.4,科目全,可伽急 +4 ELBZ6WwEaRj6 2026-11-04 8/400 2026-11-05 20:45 by qL1cU46YIdK0
[硕博家园] 售一区SCI文章T0P,我:8O.551.O54,科目全,可十急 +4 ELBZ6WwEaRj6 2026-11-04 7/350 2026-11-05 20:43 by qL1cU46YIdK0
[公派出国] 售SCI-T0P文章,我:8O.5.5.1.O.54,科目齐全,可+急 +3 ELBZ6WwEaRj6 2026-11-04 8/400 2026-11-05 20:43 by qL1cU46YIdK0
[考研] 售SCI一区文章,我:8.O.551.O.5.4,科目全,可伽急 +3 ELBZ6WwEaRj6 2026-11-04 8/400 2026-11-05 20:43 by qL1cU46YIdK0
[硕博家园] 售SCI一区文章,我:8.O.55.1.O.54,科目齐全,可伽急 +3 ELBZ6WwEaRj6 2026-11-04 9/450 2026-11-05 20:34 by qL1cU46YIdK0
[公派出国] 售SCI一区文章,我:8.O.551.O.5.4,科目全,可伽急 +4 ttH3f1oA6ace 2026-11-04 8/400 2026-11-05 20:23 by qL1cU46YIdK0
[公派出国] 售SCI一区T0P文章,我:8.O.55.1.O54,科目全,可伽急 +5 kvewwGtKktqB 2026-11-04 12/600 2026-11-05 20:21 by qL1cU46YIdK0
[硕博家园] 售SCI一区T0P文章,我:8.O55.1.O.54,科目全,可十急 +4 hNpp4Y2CGvSC 2026-11-04 11/550 2026-11-05 20:21 by qL1cU46YIdK0
[公派出国] 售SCI一区T0P文章,我:8O.55.1.O.5.4,科目齐全,可+急 +5 ncTVBM6AOJHF 2026-11-04 12/600 2026-11-05 20:10 by qL1cU46YIdK0
[考博] 售SCI一区T0P文章,我:8.O.55.1.O.5.4,科目全,可+急 +6 ncTVBM6AOJHF 2026-11-04 9/450 2026-11-05 20:10 by qL1cU46YIdK0
[考博] 售一区SCI文章T0P,我:8O.551.O54,科目全,可十急 +3 ttH3f1oA6ace 2026-11-04 6/300 2026-11-05 18:00 by My20PsvIExYO
[考研] 售SCI-T0P文章,我:8O.5.5.1.O.54,科目齐全,可+急 +3 ttH3f1oA6ace 2026-11-04 11/550 2026-11-05 17:58 by My20PsvIExYO
[考研] 售SCI一区T0P文章,我:8O.55.1.O.5.4,科目齐全,可+急 +5 ncTVBM6AOJHF 2026-11-04 10/500 2026-11-05 17:36 by My20PsvIExYO
[硕博家园] 售SCI一区文章,我:8.O.551.O.5.4,科目全,可伽急 +4 ncTVBM6AOJHF 2026-11-04 9/450 2026-11-05 17:26 by My20PsvIExYO
[硕博家园] 售SCI一区文章,我:8O5.5.1.O5.4,科目全,可伽急 +3 ncTVBM6AOJHF 2026-11-04 9/450 2026-11-05 17:25 by My20PsvIExYO
[考博] 售SCI一区文章,我:8.O.55.1.O.54,科目齐全,可伽急 +6 ncTVBM6AOJHF 2026-11-04 16/800 2026-11-05 17:25 by My20PsvIExYO
[考研] 售一区SCI文章T0P,我:8O.551.O54,科目全,可十急 +4 ncTVBM6AOJHF 2026-11-04 13/650 2026-11-05 17:15 by My20PsvIExYO
[公派出国] 售SCI-T0P文章,我:8O.5.5.1.O.54,科目齐全,可+急 +5 ncTVBM6AOJHF 2026-11-04 9/450 2026-11-05 15:13 by My20PsvIExYO
[硕博家园] 售SCI一区T0P文章,我:8O.55.1.O.54,科目全,可伽急 +4 ncTVBM6AOJHF 2026-11-04 10/500 2026-11-05 10:39 by PYZmcHr5kSaq
[基金申请] 各位基础老师们,看完中央《从严治党》的决定后,对学术界更加有信心了吧 +5 iaeyuan 2026-11-01 5/250 2026-11-04 14:43 by jinchengzi
信息提示
请填处理意见