列出帖子

该操作将允许你查看该会员所有的帖子,注意你只能看到你有权限看到的板块的帖子。


显示所有帖子 - 380025887

页: [1]
1
VX同QQ

2
广告和有偿招募 / 有偿做各种cpu程序转cuda的加速优化
« 于: 六月 12, 2019, 10:25:41 am »
如题,如有需要加QQ380025887,验证信息填cuda

3
CUDA / 稀疏矩阵问题
« 于: 八月 21, 2018, 05:39:22 pm »
我想请问下AX=B求解X的问题,A是一个大型正定稀疏矩阵240000*240000,我按照cudaSample中的例子(cuSolverSp_LinearSolver)使用cusolver.library中的cholesky分解法来解大概需要6秒左右,而matlab的mldivide函数只要1s,按理说cuda应该更快才是呀,我想知道是我使用方法不对还是cuda本身不适合解这种方程组。另外使用MKL的话大概是1.5S


4
CUDA / 这段matlab代码转为cuda该怎么优化
« 于: 七月 02, 2018, 05:15:14 pm »
程序代码: [选择]
LPdist2 = LxPx.^2 + LyPy.^2 + LzPz.^2;
LPdist = sqrt(LPdist2);
LxPx = LxPx ./ LPdist; LyPy = LyPy ./ LPdist; LzPz = LzPz ./ LPdist;
imgSet2Ddist = imgSet2D .* LPdist2 / max(LPdist2(:));
这是一段matlab代码,所有标识符都是大小为22*2304000的矩阵,由于所有运算都是点对点的运算,一开始我的想法是每一个线程计算一个像素点,但是到最后一行代码的时候发现必须先求取LPdist2矩阵中的最大值,无法继续点乘点除,我想问下有什么解决方法。我目前只能想到先求出LPdist2矩阵中所有点的值再求出最大值,但是感觉效率会很低。

5
CUDA /
« 于: 六月 29, 2018, 02:38:49 pm »
楼主直接使用,例如1221 * 512的形状启动即可, 然后:
int tid = blockIdx.x * blockDim.x + threadIdx ...

非常感谢版主, 我先去阅读一下profiler的手册吧,刚刚用profiler只是知道个运行时间,不是很懂怎么看瓶颈。小白才刚学一周

6
CUDA /
« 于: 六月 29, 2018, 10:07:06 am »
形状配置无固定标准,也没有通用的公式的。

你只能从多种可能的形状中选择一个对你特定的代码效果最好 ...

你好,我的代码现在展开后优化成这个样子,其中gpu_orig为625*1000*602的内存块,num_size是602,row是625,colum是1000,FinalB1是最终求得的625*1000的数据,对gpu_orig读取两次,总共是3G,显卡是192G的带宽,理论值应该是3G/192G=15MS左右(这么算没错吧),我这个实际下来是150MS,我想问下还有哪些地方可以优化。代码如下:
 
程序代码: [选择]
__global__ void cal_kernel(float*gpu_orig, int num_size, int row, int colum,float *FinalB1)
{
        const int idx = (blockIdx.x*blockDim.x) + threadIdx.x;
        const int idy = (blockIdx.y*blockDim.y) + threadIdx.y;
        const int thread_idx = ((gridDim.x*blockDim.x)*idy) + idx;
        if (thread_idx >= row*colum) return;
        float temp_max = 0;
        float temp_min = 0,temp_num;
        for (int j = 0; j < num_size; j++)//得到每个矩阵点在602个矩阵中的最大值和最小值
        {
                temp_num = gpu_orig[j*colum*row + thread_idx];
                if (j == 0)
                {
                        temp_max = temp_num;
                        temp_min = temp_num;
                }
                else
                {
                        if (temp_num > temp_max)
                        {
                                temp_max = temp_num;
                                continue;
                        }
                        if (temp_num < temp_min)
                        {
                                temp_min = temp_num;
                                continue;
                        }
                }
        }
        temp_max = temp_max - temp_min;
        float X0 = 0, X1 = 0, X2 = 0, X3 = 0, X4 = 0, Y1 = 0, Y2 = 0, Y3 = 0;
        for (int j = 0; j < num_size; j++)
        {
                temp_num = gpu_orig[j*colum*row + thread_idx];
                float temp_range1;
                float temp_Ind2sharps;
                temp_num = temp_num - temp_min;
                if (temp_num < (temp_max *0.25))
                {
                        temp_range1 = 0;
                        temp_Ind2sharps = 0;
                        temp_num = 0;
                }
                else
                {
                        temp_range1 = 1;
                        temp_Ind2sharps = j + 1;
                }
                X0 += temp_range1;
                X1 += temp_Ind2sharps;
                X2 += temp_Ind2sharps * temp_Ind2sharps;
                X3 += temp_Ind2sharps * temp_Ind2sharps * temp_Ind2sharps;
                X4 += temp_Ind2sharps * temp_Ind2sharps * temp_Ind2sharps * temp_Ind2sharps;
                Y1 += temp_Ind2sharps * temp_Ind2sharps * temp_num;
                Y2 += temp_Ind2sharps * temp_num;
                Y3 += temp_num;
        }
                float X22 = 0, X23 = 0, X33 = 0, XX3 = 0, XX2 = 0, C1 = 0, B1 = 0, A1 = 0, conf = 0;
                X22 = X2 - X3*X3 / X4;
                X23 = X1 - X2*X3 / X4;
                Y2 = Y2 - Y1*X3 / X4;
                X33 = X0 - X2*X2 / X4;
                Y3 = Y3 - Y1*X2 / X4;
                X33 = X33 - X23*X23 / X22;
                Y3 = Y3 - Y2*X23 / X22;
                XX3 = Y3 / X33;
                XX2 = (Y2 - X23*XX3) / X22;
                C1 = (Y1 - X2*XX3 - X3*XX2) / X4;
                B1 = -XX2 / (C1 * 2);
                A1 = -(XX2*XX2 / (4 * C1)) + XX3;
                conf = A1 / C1;
                if (B1 > num_size || B1 < 1 || C1>0 || A1 < 0 || conf<-2500 || conf>-10)
                {
                        FinalB1[thread_idx] = 0;
                }
                else
                {
                        FinalB1[thread_idx] = B1;
                }
}



7
资源下载 /
« 于: 六月 28, 2018, 09:43:17 am »
感谢楼主提供资源

8
CUDA /
« 于: 六月 28, 2018, 09:18:33 am »
形状配置无固定标准,也没有通用的公式的。

你只能从多种可能的形状中选择一个对你特定的代码效果最好 ...

感谢[名词6]

9
CUDA /
« 于: 六月 27, 2018, 04:34:09 pm »
你的代码是常规的已经展开了。
你所看到的for循环(除了内层的那个j, 这里应该是循环多个图层(根据你的上 ...

      您好,我按照您的思路改了下代码,看起来是优化了不少,本来以为速度会优化很多,但是用cudaEvent来记录运行时间发现速度上几乎没有提升,以前我那种写法是看了一个博客上面介绍的对内存访问为连续的方式,不知道速度没有得到提升是否跟这个有关。      博客地址为:https://blog.csdn.net/sunmc1204953974/article/details/51064302。      更改后的代码如下:
 
程序代码: [选择]
__global__ void get_max_min(float*gpu_orig, char *gpurange1, int num_size, int row, int colum, float *gpu_Ind2sharps)
{
        const int idx = (blockIdx.x*blockDim.x) + threadIdx.x;
        const int idy = (blockIdx.y*blockDim.y) + threadIdx.y;
        const int thread_idx = ((gridDim.x*blockDim.x)*idy) + idx;
        float temp_max = 0;
        float temp_min = 0;
        if (thread_idx >= row*colum)   return;
        for (int j = 0; j < num_size; j++)
        {
                if (j == 0)
                {
                        temp_max = gpu_orig[j*colum*row + thread_idx];
                        temp_min = gpu_orig[j*colum*row + thread_idx];
                }
                else
                {
                        if (gpu_orig[j*colum*row + thread_idx] > temp_max)    temp_max = gpu_orig[j*colum*row + thread_idx];
                        if (gpu_orig[j*colum*row + thread_idx] < temp_min)     temp_min = gpu_orig[j*colum*row + thread_idx];
                }

        }
        temp_max = temp_max - temp_min;
        for (int j = 0; j < num_size; j++)
        {
                gpu_orig[j * row * colum + thread_idx] = gpu_orig[j * row * colum + thread_idx] - temp_min;
                if (gpu_orig[j * row * colum + thread_idx] < (temp_max *0.25))
                {
                        gpurange1[j * row * colum + thread_idx] = 0;

                        gpu_orig[j * row * colum + thread_idx] = 0;
                }
                else
                {
                        gpurange1[j * row * colum + thread_idx] = 1;
                }
                gpu_Ind2sharps[j * row * colum + thread_idx] = (j + 1)*gpurange1[j * row * colum + thread_idx];
        }
}





10
CUDA /
« 于: 六月 27, 2018, 02:44:50 pm »
你的代码是常规的已经展开了。
你所看到的for循环(除了内层的那个j, 这里应该是循环多个图层(根据你的上 ...

感谢版主提供的思路,还有个问题就是这个block的数目以及使用2维三维还有每个bolck中thread的数量有什么讲究吗,我只知道每个block中thread应该是32的倍数。就比如说我设置每个线程块中线程数量为192按照您的说法应该是使用625*1000/192约等于32556个线程块,这些线程块和线程我该怎么设置比较合理。

11
CUDA /
« 于: 六月 27, 2018, 10:26:28 am »
这能删除吗。。。为什么发出来的和编辑的不一样啊。就是gpu_max和gpu_min后面缺一个,
比如gpu_max = temp_max;gpu_min = temp_min;

12
CUDA /
« 于: 六月 27, 2018, 10:23:27 am »
楼主,

你的原始问题没有看懂。


又乱了,再发下。
程序代码: [选择]
__global__ void get_max_min(float*gpu_orig,char *gpurange1,float*gpu_max, float *gpu_min, int num_size, int row, int colum,float *gpu_Ind2sharps)
{
        const int tid = threadIdx.x;
        const int bid = blockIdx.x;
        float temp_max = 0;
        float temp_min = 0;
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        if (j == 0)
                        {
                                temp_max = gpu_orig[j*colum*row + i];
                                temp_min = gpu_orig[j*colum*row + i];
                        }
                        else
                        {
                                if (gpu_orig[j*colum*row + i] > temp_max) temp_max = gpu_orig[j*colum*row + i];
                                if (gpu_orig[j*colum*row + i] < temp_min) temp_min = gpu_orig[j*colum*row + i];
                        }
                       
                }
                gpu_max = temp_max;
                gpu_min = temp_min;
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                gpu_max = gpu_max - gpu_min;
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        gpu_orig[j * row * colum + i] = gpu_orig[j * row * colum + i] - gpu_min;
                }
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        if (gpu_orig[j * row * colum + i] < (gpu_max *0.25))
                        {
                                gpurange1[j * row * colum + i] = 0;


                                gpu_orig[j * row * colum + i] = 0;
                        }
                        else
                        {
                                gpurange1[j * row * colum + i] = 1;
                        }
                }
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        gpu_Ind2sharps[j * row * colum + i] = (j + 1)*gpurange1[j * row * colum + i];
                }
        }
}



13
CUDA /
« 于: 六月 27, 2018, 10:13:30 am »
楼主,

你的原始问题没有看懂。


上个回复代码乱了,重新回复下代码
程序代码: [选择]
__global__ void get_max_min(float*gpu_orig,char *gpurange1,float*gpu_max, float *gpu_min, int num_size, int row, int colum,float *gpu_Ind2sharps)
{
        const int tid = threadIdx.x;
        const int bid = blockIdx.x;
        float temp_max = 0;
        float temp_min = 0;
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        if (j == 0)
                        {
                                temp_max = gpu_orig[j*colum*row + i];
                                temp_min = gpu_orig[j*colum*row + i];
                        }
                        else
                        {
                                if (gpu_orig[j*colum*row + i] > temp_max) temp_max = gpu_orig[j*colum*row + i];
                                if (gpu_orig[j*colum*row + i] < temp_min) temp_min = gpu_orig[j*colum*row + i];
                        }
                       
                }
                gpu_max = temp_max;
                gpu_min = temp_min;
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                gpu_max = gpu_max - gpu_min;
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        gpu_orig[j * row * colum + i] = gpu_orig[j * row * colum + i] - gpu_min;
                }
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        if (gpu_orig[j * row * colum + i] < (gpu_max *0.25))
                        {
                                gpurange1[j * row * colum + i] = 0;

                                gpu_orig[j * row * colum + i] = 0;
                        }
                        else
                        {
                                gpurange1[j * row * colum + i] = 1;
                        }
                }
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        gpu_Ind2sharps[j * row * colum + i] = (j + 1)*gpurange1[j * row * colum + i];
                }
        }
}



14
CUDA /
« 于: 六月 27, 2018, 10:06:27 am »
楼主,

你的原始问题没有看懂。


感谢版主的回复,目前我的单个嵌套是625*1000*602三层for循环嵌套(602张像素大小为625*1000的图片),然后同时也有几个相同的嵌套按顺序并行,我现在想问的是这些并行的嵌套是写在一个kernel中比较好还是多个kernel比较好。附上我的一个kernel代码,我将625和1000合并了,按我目前这样写也可以运行,我使用的CUDA版本是9.2.88+VS2015。还有我的for循环还有什么比较好的优化方法吗
程序代码: [选择]
__global__ void get_max_min(float*gpu_orig,char *gpurange1,float*gpu_max, float *gpu_min, int num_size, int row, int colum,float *gpu_Ind2sharps)
{
        const int tid = threadIdx.x;
        const int bid = blockIdx.x;
        float temp_max = 0;
        float temp_min = 0;
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        if (j == 0)
                        {
                                temp_max = gpu_orig[j*colum*row + i];
                                temp_min = gpu_orig[j*colum*row + i];
                        }
                        else
                        {
                                if (gpu_orig[j*colum*row + i] > temp_max) temp_max = gpu_orig[j*colum*row + i];
                                if (gpu_orig[j*colum*row + i] < temp_min) temp_min = gpu_orig[j*colum*row + i];
                        }
                       
                }
                gpu_max = temp_max;
                gpu_min = temp_min;        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                gpu_max = gpu_max - gpu_min;
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        gpu_orig[j * row * colum + i] = gpu_orig[j * row * colum + i] - gpu_min;
                }
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        if (gpu_orig[j * row * colum + i] < (gpu_max *0.25))
                        {
                                gpurange1[j * row * colum + i] = 0;


                                gpu_orig[j * row * colum + i] = 0;
                        }
                        else
                        {
                                gpurange1[j * row * colum + i] = 1;
                        }
                }
        }
        for (int i = bid*THREAD_NUM + tid; i < row*colum; i += BLOCK_NUM * THREAD_NUM)
        {
                for (int j = 0; j < num_size; j++)
                {
                        gpu_Ind2sharps[j * row * colum + i] = (j + 1)*gpurange1[j * row * colum + i];
                }
        }
}



15
CUDA / kernel函数大小问题
« 于: 六月 26, 2018, 02:53:11 pm »
现在我需要处理多个for循环,是将这些for循环写入一个核函数中比较好还是写多个核函数分别调用比较好。本人刚刚开始学CUDA,希望各位[名词6]不要见笑

页: [1]