找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 49|回复: 1

CUDA矩阵乘法示例代码,调大矩阵维度后崩溃

[复制链接]
发表于 7 天前 | 显示全部楼层 |阅读模式
ESC4000G3
大家好,最近在学习CUDA,自己手敲了矩阵乘法的一个Demo程序,发现在矩阵维度比较小的时候都很顺利,但调大矩阵维数后显卡就崩溃了。但是算了一下,4096*4096的矩阵也就几十兆,也是分配在Global Memory下的,显卡用的是GT650,应该不会出问题啊,求各位大佬帮忙看一下。
代码如下:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include<time.h>
using namespace std;
#define TILE_WIDTH 32


__global__ void MatrixMulKernleGlobal(int m, int n, int k, float *A,float  *B, float *C)
{
        int Row=blockIdx.y*blockDim.y+threadIdx.y;
        int Col=blockIdx.x*blockDim.x+threadIdx.x;

        if((Row<m) && (Col<k))
        {
                float Cvalue=0.0;
                for(int i=0;i<n;++i)
                     Cvalue+=A[Row*n+i]*B[Col+i*k];
                C[Row*k+Col]=Cvalue;
        }
}

int main()
{
        int height = (1<<12);
        int width = (1<<12);
        int length = (1<<12);

        cout << height << ' ' << width << ' ' << length << ' ' << endl;

        // allocate memory
        float *A, *B, *C;
        A = (float*)malloc(height*width*sizeof(float));
        B = (float*)malloc(width*length*sizeof(float));
        C = (float*)malloc(height*length*sizeof(float));

        // initialize
        for(int i=0; i<height*width; ++i)
                A[i] = 1.0;
        for(int i=0; i<width*length; ++i)
                B[i] = 2.0;

        //分配显存空间
        int size = sizeof(float);
        float *d_a;
        float *d_b;
        float *d_c;
        cudaMalloc((void**)&d_a,height*width*size);
        cudaMalloc((void**)&d_b,width*length*size);
        cudaMalloc((void**)&d_c,height*length*size);

        //把数据从Host传到Device
        cudaMemcpy(d_a, A, size*height*width, cudaMemcpyHostToDevice);
        cudaMemcpy(d_b, B, size*width*length, cudaMemcpyHostToDevice);

        //分配网格结构
        dim3 dimGrid((length-1)/TILE_WIDTH+1,(height-1)/TILE_WIDTH+1);        //向上取整
        dim3 dimBlock(TILE_WIDTH,TILE_WIDTH);

        cout << dimGrid.x  << ' ' << dimGrid.y << ' ' << dimGrid.z << endl;
        cout << dimBlock.x << ' ' << dimBlock.y << ' ' << dimBlock.z << endl;

        //调用内核函数
        MatrixMulKernleGlobal<<<dimGrid,dimBlock>>>(height,width,length,d_a,d_b,d_c);

        //将结果传回到主机端
        cudaMemcpy(C, d_c, size*height*length, cudaMemcpyDeviceToHost);

        //输出结果
        for (int i=0;i<8;i++)
        {
                cout<<C[i]<<endl;
        }

        //释放显存空间
        cudaFree(d_a);
        cudaFree(d_b);
        cudaFree(d_c);

        return 0;
}

回复

使用道具 举报

发表于 7 天前 | 显示全部楼层
Jetson TX2
本帖最后由 屠戮人神 于 2018-7-13 11:07 编辑

楼主你好,

你这种实现的矩阵乘法是非常native的, 访存极度低效。
例如你这里的Cvalue+=A[Row*n+i]*B[Col+i*k];
这种代码规模一大必然超时,加上你的卡太渣(GT650,目前价格80元淘宝),更容易凸显这个问题。

也可以看下网上的例子, 随手找了一个(有图):http://www.compsci.hunter.cuny.e ... UDA_matrix_mult.pdf

建议的解决方案:
(1)改成cublas直接完成。
(2)或者考虑使用shared memory的版本(CUDA手册上就有)
(3)或者购买更加强大的卡(这样可以在更大的规模才会挂---依然会挂的。只是挂的时机)。

不排除你的代码还有更多问题, 但这个是主要问题。



回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表