# CUDA之矩阵乘法——复数

### 代码

__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;
int row = threadIdx.y;
int col = threadIdx.x;
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);
}

• 本文已收录于以下专栏：

## CUDA快速傅里叶变换 - cuFFT

CUDA为开发人员提供了多种库，每一类库针对某一特定领域的应用，CUFFT库则是CUDA中专门用于进行傅里叶变换的函数库，这一系列的文章是博主近一段时间对cuFFT库的学习总结，主要内容是文档的译文，...
• Augusdi
• 2013年10月10日 10:20
• 11340

## CUDA复数类的定义

typedef float2 cuFloatComplex; __host__ __device__ static __inline__ float cuCrealf (cuFloatComplex ...
• crouse246
• 2015年05月25日 11:44
• 382

## 终于成功配置numba cuda

• u013975830
• 2017年12月16日 23:21
• 329

## Numba CUDA 初试

• m0_38009702
• 2017年10月04日 16:40
• 259

## CUDA 矩阵转置

• chinacodec
• 2009年03月02日 17:30
• 5525

## CUDA之矩阵乘法——复数

• Sumujingling
• 2016年05月30日 16:55
• 544

## 实数及复数矩阵加法并行CUDA

• 2014年08月28日 00:54
• 19.8MB
• 下载

## CUDA: 矩阵乘法优化

• Augusdi
• 2013年10月12日 10:39
• 9837

## 矩阵乘法——CUDA 优化记录

CUDA 是 NVIDIA 的 GPGPU 模型，它使用 C 语言为基础，可以直接以大多数人熟悉的 C 语言，写出在显示芯片上执行的程序，而不需要去学习特定的显示芯片的指令或是特殊的结构。” 编...
• Openking
• 2015年01月09日 10:01
• 1815

## 复数四则运算(利用C++的复数库实现复数的各种数学运算)

• 2010年04月22日 20:16
• 1KB
• 下载

举报原因： 您举报文章：CUDA之矩阵乘法——复数 色情 政治 抄袭 广告 招聘 骂人 其他 (最多只允许输入30个字)