cuda编程------矩阵乘法

cuda + VS 2010 安装：http://www.cnblogs.com/xing901022/archive/2013/08/09/3248469.html

本文主要介绍如何使用CUDA并行计算矩阵乘法：

//头文件
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void MatMulKernel(float* A, float* B, float* C,dim3 dimsA, dim3 dimsB)//[1]
{
// Each thread computes one element of C
// by accumulating results into Cvalue

float Cvalue = 0;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
for (int e = 0; e < dimsA.y; ++e)
Cvalue += A[row * dimsA.y + e]* B[e * dimsB.y + col];
C[row * dimsB.y  + col] = Cvalue;
}

void MatrixMultiplication_CUDA(const float* A,const float* B,float* C)
{
dim3 dimsA(4,3);// the size of matrix A which you have to modify
dim3 dimsB(3,3);// the size of matirx B which you have to modify

//copy memory from host to devices
unsigned int size_A = dimsA.x * dimsA.y;
unsigned int mem_size_A = sizeof(float) * size_A;
float *d_A ;
cudaMalloc(&d_A,mem_size_A);
cudaMemcpy(d_A,A,mem_size_A,cudaMemcpyHostToDevice);
float *d_B;
unsigned int size_B = dimsB.x * dimsB.y;
unsigned int mem_size_B = sizeof(float) * size_B;
cudaMalloc(&d_B,mem_size_B);
cudaMemcpy(d_B,B,mem_size_B,cudaMemcpyHostToDevice);
unsigned int mem_size_C = sizeof(float)* dimsA.x*dimsB.y;
float *d_C;
cudaMalloc(&d_C,mem_size_C);

//dimBlock represents the threads'size within block which you have to modify[2]
dim3 dimBlock(3,2);
dim3 dimGrid(dimsB.y/dimBlock.x,dimsA.x/dimBlock.y);//[3]

MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C, dimsA,  dimsB);
// Read C from device memory
cudaMemcpy(C, d_C, mem_size_C,
cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);

}
int main()
{
float A[] = {6,2,3,
8,3,5,
7,2,4,
8.3,2,5};
float B[] ={1,2,3,
4,5,6,
7,8,9};
float*C = new float[12];
MatrixMultiplication_CUDA(A,B,C);
for(int i =0;i<12;i++)
printf("%f ",C[i]);
getchar();
return 0;
}

[1] __gloabal__函数

A kernel is defined using the __global__ declaration specifier and the number of CUDA threads that execute that kernel for a given kernel call is specified using a new

<<<...>>> execution configuration syntax

使用__global__声明的函数即是CUDA 的kernel函数，也就是需要并行计算的部分，线程数目则是调用时由<<<>>>指定。

[2] dimBlock 指定的是一个block中的线程数目，dimGrid指定的是一个grid中block的数目。一个grid由多个block组成，一个block内有多个线程，block之间不能交互，但一个block里面的线程可以通过shared memory进行交互。

这个可以自己指定，但必须满足条件 dimBlock.x< B.y dimBlock.y<A.x。否则[3]指定的block就有可能为0了 。而且一个block最多有512个线程

例如main函数中A.x = 4 我设置的dimBlock.y为2 ，2<4

[3]可能疑惑X,Y是否弄反了，看gpu里面排列表就知道了。纵列和我们的数组排列就是不一样。