# CUDA之矩阵乘法——复数

205人阅读 评论(0)

### 代码

__device__ float GetReal(const Matrix A, int row, int col) {

return A.real[row * A.stride + col];
}

__device__ float GetImag(const Matrix A, int row, int col) {
return A.imag[row * A.stride + col];
}

__device__ void SetElement(Matrix A, int row, int col, float valueR, float valueI) {
A.real[row * A.stride + col] = valueR;
A.imag[row * A.stride + col] = valueI;
}

__device__ Matrix GetSubMatrix(Matrix A, int row, int col) {
Matrix Asub;
Asub.width = BLOCK_SIZE;
Asub.height = BLOCK_SIZE;
Asub.stride = A.stride;
Asub.real = &A.real[A.stride * BLOCK_SIZE * row+ BLOCK_SIZE * col];
Asub.imag = &A.imag[A.stride * BLOCK_SIZE * row+ BLOCK_SIZE * col];
return Asub;
}

__global__ void CMatMulKernel(Matrix A, Matrix B, Matrix C) {
int blockRow = blockIdx.y;
int blockCol = blockIdx.x;
Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
float CvalueR = 0;
float CvalueI = 0;
for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
Matrix Asub = GetSubMatrix(A, blockRow, m);
Matrix Bsub = GetSubMatrix(B, m, blockCol);
__shared__ float AsR[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float AsI[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float BsR[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float BsI[BLOCK_SIZE][BLOCK_SIZE];
AsR[row][col] = GetReal(Asub, row, col);
AsI[row][col] = GetImag(Asub, row, col);
BsR[row][col] = GetReal(Bsub, row, col);
BsI[row][col] = GetImag(Bsub, row, col);
for (int e = 0; e < BLOCK_SIZE; ++e)
{
CvalueR += AsR[row][e] * BsR[e][col]-AsI[row][e]*BsI[e][col];
CvalueI += AsR[row][e] * BsI[e][col]+AsI[row][e]*BsR[e][col];
}
}
SetElement(Csub, row, col, CvalueR,CvalueI);
}

void CMatMul(const Matrix A, const Matrix B, Matrix C) {
Matrix d_A;
d_A.width = d_A.stride = A.width;
d_A.height = A.height;
size_t size = A.width * A.height * sizeof(float);
cudaMalloc((void**)&d_A.real, size);
cudaMalloc((void**)&d_A.imag, size);
cudaMemcpy(d_A.real, A.real, size,
cudaMemcpyHostToDevice);
cudaMemcpy(d_A.imag, A.imag, size,
cudaMemcpyHostToDevice);

Matrix d_B;
d_B.width = d_B.stride = B.width;
d_B.height = B.height;
size = B.width * B.height * sizeof(float);
cudaMalloc((void**)&d_B.real, size);
cudaMalloc((void**)&d_B.imag, size);
cudaMemcpy(d_B.real, B.real, size,
cudaMemcpyHostToDevice);
cudaMemcpy(d_B.imag, B.imag, size,
cudaMemcpyHostToDevice);

Matrix d_C;
d_C.width = d_C.stride = C.width;
d_C.height = C.height;
size = C.width * C.height * sizeof(float);
cudaMalloc((void**)&d_C.real, size);
cudaMalloc((void**)&d_C.imag, size);

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
CMatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
cudaMemcpy(C.real, d_C.real, size,
cudaMemcpyDeviceToHost);
cudaMemcpy(C.imag, d_C.imag, size,
cudaMemcpyDeviceToHost);

cudaFree(d_A.real);
cudaFree(d_A.imag);
cudaFree(d_B.real);
cudaFree(d_B.imag);
cudaFree(d_C.real);
cudaFree(d_C.imag);
}

## 为什么这么久过去了我依然不能上传图片。

0
0

* 以上用户言论只代表其个人观点，不代表CSDN网站的观点或立场
个人资料
• 访问：11465次
• 积分：294
• 等级：
• 排名：千里之外
• 原创：18篇
• 转载：1篇
• 译文：0篇
• 评论：1条
文章分类
评论排行
最新评论