CUDA 高性能并行计算入门

CUDA 高性能并行计算入门

概述

GPU硬件介绍

下面具体介绍硬件架构

__Host__

__Global__

__Shared__

__Device__

__Constant__

程序性能分析

在Ubuntu上配置CUDA

//运行run文件
sudo sh cuda_8.xx.xx_linux.run
//编辑环境变量
sudo gedit /etc/profile
//在末尾加入以下环境变量
export PATH=/usr/local/cuda-8.0/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda-8.0/lib64$LD_LIBRARY_PATH

nvcc

额外的程序环境

小试牛刀

#include <stdio.h>
#include <iostream>
#include <cuda_runtime.h>
using namespace std;

int size = 5;
float *a = (float*)malloc(sizeof(float)*size*size);
for (int i = 0; i < size; i++) {
for (int j = 0; j < size; j++) {
a[i*size+j] = 1.0f;
}
}

float *a_cuda;
cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);

cudaMemcpy(a_cuda,a,sizeof(float)*size*size,cudaMemcpyHostToDevice);

Function<<<griddim,blockdim,extern shared memory,GPU stream>>>(param...);

__global__

__global__ void addone(float *a) {
int tix = threadIdx.x;
int tiy = threadIdx.y;
int bdx = blockDim.x;
int bdy = blockDim.y;
a[tix*bdy+tiy] += 1;
}

dim3 grid(1, 1, 1), block(5, 5, 1);
addone<<<grid,block>>>(a_cuda);

cudaMemcpy(a,a_cuda,sizeof(float)*size*size,cudaMemcpyDeviceToHost);

void MatrixPrint(float *mat, int rows, int cols) {
for (int i = 0; i < rows; i++) {
for (int j = 0; j < cols; j++) {
cout << setw(2) << mat[i*cols+j] << " ";
}
cout << endl;
}
cout << endl;
}

#include <iomanip>

#include <stdio.h>
#include <iostream>
#include <iomanip>
#include <cuda_runtime.h>
using namespace std;

void MatrixPrint(float *mat, int rows, int cols) {
for (int i = 0; i < rows; i++) {
for (int j = 0; j < cols; j++) {
cout << setw(2) << mat[i*cols+j] << " ";
}
cout << endl;
}
cout << endl;
}

__global__ void addone(float *a) {
int tix = threadIdx.x;
int tiy = threadIdx.y;
int bdx = blockDim.x;
int bdy = blockDim.y;
a[tix*bdy+tiy] += 1;
}
int main()
{
int size = 5;
float *a = (float*)malloc(sizeof(float)*size*size);
for (int i = 0; i < size; i++) {
for (int j = 0; j < size; j++) {
a[i*size+j] = 1.0f;
}
}
MatrixPrint(a,size,size);
float *a_cuda;
cudaMalloc((void**)&a_cuda,sizeof(float)*size*size);
cudaMemcpy(a_cuda,a,sizeof(float)*size*size,cudaMemcpyHostToDevice);

dim3 grid(1, 1, 1), block(5, 5, 1);
addone<<<grid,block>>>(a_cuda);
cudaMemcpy(a,a_cuda,sizeof(float)*size*size,cudaMemcpyDeviceToHost);
MatrixPrint(a,size,size);
return 0;
}

nvcc -gencode=arch=compute_50,code=\"sm_50,compute_50\" -o basic basic.cu

./basic

进阶之路

矩阵相乘（第“3的0次方”讲）

void MatrixMul_host(float *a, int a_rows, int a_cols, float *b, int b_rows, int b_cols, float *c) {
for (int i = 0; i < a_rows; i++) {
for (int j = 0; j < b_cols; j++) {
float t = 0;
for (int k = 0; k < b_rows; k++) {
t += a[i*a_cols+k]*b[k*b_cols+j];
}
c[i*b_cols+j] = t;
}
}
}

void MatrixRandBin(float *mat, int rows, int cols) {
for (int i = 0; i < rows; i++) {
for (int j = 0; j < cols; j++) {
if ((float)rand()/RAND_MAX > 0.5) {
mat[i*cols+j] = 1.0f;
}else {
mat[i*cols+j] = -1.0f;
}
}
}
}

float MatrixCompare(float *a,float *b,int rows,int cols){
float err=0;
for (int i=0;i<rows;i++){
for (int j=0;j<cols;j++){
err+=abs(a[i*cols+j]-b[i*cols+j]);
}
}
return err;
}

__global__ void MatrixMul_device(float *a, int a_rows, int a_cols, float *b, int b_rows, int b_cols, float *c) {
int tix = threadIdx.x;
int tiy = threadIdx.y;
int bix = blockIdx.x;
int biy = blockIdx.y;
int bdx = blockDim.x;
int bdy = blockDim.y;
int gdx = gridDim.x;
int gdy = gridDim.y;

for (int i = tix; i < b_cols; i += bdx) {
float sum = 0;
for (int k = 0; k < a_cols; k++) {
sum += a[bix*a_rows+k]*b[k*b_cols+i];
}
c[bix*a_cols+i] = sum;
}
}

int Matrixsize=1000;
float *a_host;
float *a_device;
float *b_host;
float *b_device;
float *result_host;
float *result_device;
float *result_cpu;
a_host = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
b_host = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
result_host = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
result_cpu = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
srand(0);
MatrixRandBin(a_host,Matrixsize,Matrixsize);
MatrixRandBin(b_host,Matrixsize,Matrixsize);
cudaMalloc((void**)&a_device,sizeof(float) *Matrixsize * Matrixsize);
cudaMalloc((void**)&b_device,sizeof(float) *Matrixsize * Matrixsize);
cudaMalloc((void**)&result_device,sizeof(float) *Matrixsize * Matrixsize);
cudaMemcpy(a_device,a_host,sizeof(float) *Matrixsize * Matrixsize,cudaMemcpyHostToDevice);
cudaMemcpy(b_device,b_host,sizeof(float) *Matrixsize * Matrixsize,cudaMemcpyHostToDevice);

cudaEvent_t start_device, stop_device;
float time_device;
cudaEventCreate(&start_device);
cudaEventCreate(&stop_device);
cudaEventRecord( start_device, 0 );
dim3 gridsize(1000,1,1);
dim3 blocksize(256,1,1);
MatrixMul_device<<<gridsize,blocksize>>>(a_device,Matrixsize,Matrixsize,b_device,Matrixsize,Matrixsize,result_device);
cudaEventRecord( stop_device, 0 );
cudaEventSynchronize( stop_device );
cudaEventElapsedTime( &time_device, start_device, stop_device );
cudaEventDestroy( start_device );
cudaEventDestroy( stop_device );
cout<<"gputime="<<time_device<<"ms"<<endl;

cudaMemcpy(result_host, result_device,sizeof(float) *Matrixsize * Matrixsize,cudaMemcpyDeviceToHost);
cudaFree(a_device);
cudaFree(b_device);
cudaFree(result_device);
clock_t start_host = clock();
MatrixMul_host(a_host,Matrixsize,Matrixsize,b_host,Matrixsize,Matrixsize,result_cpu);
cout<<"cputime="<<(double)(clock() - start_host)/1000<<"ms"<<endl;
float err=MatrixCompare(result_cpu,result_host,Matrixsize,Matrixsize);
cout<<"err in gpu and cpu = "<<err<<endl;

矩阵相乘（第“In(e的2次方)”讲）

size_t pitch_a_device, pitch_b_device, pitch_result_device;
cudaMallocPitch((void**)&a_device , &pitch_a_device , sizeof(float) *Matrixsize , Matrixsize);
cudaMallocPitch((void**)&b_device , &pitch_b_device , sizeof(float) *Matrixsize , Matrixsize);
cudaMallocPitch((void**)&result_device , &pitch_result_device , sizeof(float) *Matrixsize , Matrixsize);
cudaMemcpy2D(a_device,pitch_a_device,a_host,sizeof(float) *Matrixsize ,sizeof(float) *Matrixsize, Matrixsize,cudaMemcpyHostToDevice);
cudaMemcpy2D(b_device,pitch_b_device,b_host,sizeof(float) *Matrixsize ,sizeof(float) *Matrixsize, Matrixsize,cudaMemcpyHostToDevice);

cudaMemcpy2D(result_host,sizeof(float) *Matrixsize, result_device,pitch_result_device,sizeof(float) *Matrixsize , Matrixsize,cudaMemcpyDeviceToHost);

__global__ void MatrixMul_device(float *a,int a_pitch,int a_rows,int a_cols,float *b,int b_pitch,int b_rows,int b_cols,float *c,int c_pitch){
int tix=threadIdx.x;
int tiy=threadIdx.y;
int bix=blockIdx.x;
int biy=blockIdx.y;
int bdx=blockDim.x;
int bdy=blockDim.y;
int gdx=gridDim.x;
int gdy=gridDim.y;
int i,k;
for(i=tix;i<b_cols;i+=bdx){
float sum=0;
for(k=0;k<a_cols;k++){
sum+=a[bix*a_pitch+k]*b[k*b_pitch+i];
}
c[bix*a_pitch+i]=sum;
}
}

MatrixMul_device<<<gridsize,blocksize>>>(a_device,pitch_a_device/sizeof(float),Matrixsize,Matrixsize,b_device,pitch_b_device/sizeof(float),Matrixsize,Matrixsize,result_device,pitch_result_device/sizeof(float));

__shared__ float a_shared[1000];
for(int j=tix;j<a_cols;j+=bdx){
a_shared[j]=a[bix*a_pitch+j];
}
__syncthreads();
int i,k;
for(i=tix;i<b_cols;i+=bdx){
float sum=0;
for(k=0;k<a_cols;k++){
sum+=a_shared[k]*b[k*b_pitch+i];
}
c[bix*a_pitch+i]=sum;
}

__shared__

__syncthreads()

矩阵相乘（第“2×(最小自然数+2)”讲）

int Matrixsize=1000;
float *a_host;
float *b_host;
float *result_host;
a_host = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
b_host = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
result_host = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
srand(0);
MatrixRandBin(a_host,Matrixsize,Matrixsize);
MatrixRandBin(b_host,Matrixsize,Matrixsize);
// cout<<MatrixCopysize<<endl;

float *a_device;
float *b_device;
float *result_device;
cudaMalloc((void**)&a_device,sizeof(float) *Matrixsize * Matrixsize);
cudaMalloc((void**)&b_device,sizeof(float) *Matrixsize * Matrixsize);
cudaMalloc((void**)&result_device,sizeof(float) *Matrixsize * Matrixsize);
cudaMemcpy(a_device,a_host,sizeof(float) *Matrixsize * Matrixsize,cudaMemcpyHostToDevice);
cudaMemcpy(b_device,b_host,sizeof(float) *Matrixsize * Matrixsize,cudaMemcpyHostToDevice);

void MatrixMul_device(float *a,float *b,int a_rows,int a_cols,int b_cols,float *result){}

int BINSIZE=32;//size of bin2int, 32 means 0000 0000 0000 0000 0000 0000 0000 0000
int MaxBS=(a_cols-1)/BINSIZE+1;
int a_cols_Copysize=MaxBS*BINSIZE;
float *a_device;//a_rows * a_cols_Copysize
float *b_device;//a_cols_Copysize * b_cols
size_t pitch_a_device, pitch_b_device;
cudaMallocPitch((void**)&a_device , &pitch_a_device , sizeof(float) *a_cols_Copysize , a_rows);
cudaMallocPitch((void**)&b_device , &pitch_b_device , sizeof(float) *b_cols , a_cols_Copysize);

cudaMemset(a_device, 0, pitch_a_device * a_rows);
cudaMemset(b_device, 0, pitch_b_device * a_cols_Copysize);
cudaMemcpy2D(a_device,pitch_a_device,a,sizeof(float) *a_cols ,sizeof(float) *a_cols, a_rows,cudaMemcpyDeviceToDevice);
cudaMemcpy2D(b_device,pitch_b_device,b,sizeof(float) *b_cols ,sizeof(float) *b_cols, a_cols,cudaMemcpyDeviceToDevice);

int *a_device_bin;
int *b_device_bin;
size_t pitch_a_device_bin, pitch_b_device_bin;
cudaMallocPitch((void**)&a_device_bin , &pitch_a_device_bin , sizeof(int) *MaxBS , a_rows);
cudaMallocPitch((void**)&b_device_bin , &pitch_b_device_bin , sizeof(int) *b_cols , MaxBS);

//horizontal
__global__ void AMatrix2Bin(float *a,int *a_bin,int a_rows,int pitch_a,int pitch_a_bin,int MaxBS,int BINSIZE){
int tix=threadIdx.x;
// int tiy=threadIdx.y;
int bix=blockIdx.x;
// int biy=blockIdx.y;
int bdx=blockDim.x;
// int bdy=blockDim.y;
int gdx=gridDim.x;
// int gdy=gridDim.y;

int maxThreads=MaxBS*a_rows;
for(int id = bix*bdx+tix; id < maxThreads; id+=gdx*bdx) {
int rid=id/MaxBS;
int cid=id%MaxBS;

int Integer=0;
int base=1;
for (int i=0;i<BINSIZE;i++){
if (a[rid*pitch_a+(cid+1)*BINSIZE-1-i]==1.f){
Integer+=base;
}
base=base<<1;
}

a_bin[rid*pitch_a_bin+cid]=Integer;
}

}
//vetical
__global__ void BMatrix2Bin(float *b,int *b_bin,int b_cols,int pitch_b,int pitch_b_bin,int MaxBS,int BINSIZE){
int tix=threadIdx.x;
// int tiy=threadIdx.y;
int bix=blockIdx.x;
// int biy=blockIdx.y;
int bdx=blockDim.x;
// int bdy=blockDim.y;
int gdx=gridDim.x;
// int gdy=gridDim.y;

int maxThreads=MaxBS*b_cols;
for(int id = bix*bdx+tix; id < maxThreads; id+=gdx*bdx) {
int cid=id/MaxBS;
int rid=id%MaxBS;

int Integer=0;
int base=1;
for (int i=0;i<BINSIZE;i++){
if (b[((rid+1)*BINSIZE-1-i)*pitch_b+cid]==1.f){
Integer+=base;
}
base=base<<1;
}

b_bin[rid*pitch_b_bin+cid]=Integer;
}

}

>>

AMatrix2Bin<<<GS_BIN,BS_BIN>>>(a_device , a_device_bin , a_rows ,
pitch_a_device/sizeof(float) , pitch_a_device_bin/sizeof(int) , MaxBS , BINSIZE);
BMatrix2Bin<<<GS_BIN,BS_BIN>>>(b_device , b_device_bin , b_cols ,
pitch_b_device/sizeof(float) , pitch_b_device_bin/sizeof(int) , MaxBS , BINSIZE);

cudaMemcpyToSymbol（）

__device__ unsigned char __popcount_tab_device[256];//__constant__ is slower than __device__
__device__ int popcount (int x) {
return __popcount_tab_device[(x >>  0) & 0xff]
+ __popcount_tab_device[(x >>  8) & 0xff]
+ __popcount_tab_device[(x >> 16) & 0xff]
+ __popcount_tab_device[(x >> 24) & 0xff];
}

const unsigned char __popcount_tab[] = {
0,1,1,2,1,2,2,3,1,2,2,3,2,3,3,4,1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,
1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,
1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,
2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,
1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,
2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,
2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,
3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,4,5,5,6,5,6,6,7,5,6,6,7,6,7,7,8,
};
cudaMemcpyToSymbol(__popcount_tab_device, __popcount_tab, sizeof(__popcount_tab));

float *result_device;//a_rows * b_cols
size_t pitch_result_device;
cudaMallocPitch((void**)&result_device , &pitch_result_device , sizeof(float) *b_cols , a_rows);

__global__ void MatrixMulXnor(int *a,int *b,int a_rows,int a_cols,
int b_cols,float *result,int pitch_a,int pitch_b,
int pitch_result,int BINSIZE,int RealMidSize){

int tix=threadIdx.x;
// int tiy=threadIdx.y;
int bix=blockIdx.x;
// int biy=blockIdx.y;
int bdx=blockDim.x;
// int bdy=blockDim.y;
int gdx=gridDim.x;
// int gdy=gridDim.y;

int rest=(BINSIZE*a_cols-RealMidSize);

for(int j=tix;j<b_cols;j+=bdx){
// printf("i=%d ; j=%d\n",i,j);
int sum=0;
for(int k=0;k<a_cols;k++){
int bin=(a[bix*pitch_a+k]^b[k*pitch_b+j]);
int negnum=popcount(bin);
int posnum=BINSIZE-negnum;
//calculate ignores the rest of BINSIZE if the Matsize cant devided by BINSIZE ,it can cause err
//(10/00)'(01/00) should be 0000 but it is 0011,so 1+1 is trash in the result.and it can cause a_rows*b_cols times.
sum+=(posnum-negnum);
}
result[bix*pitch_result+j]=sum-rest;
}
}

cudaEvent_t start_device, stop_device;
float time_device;
cudaEventCreate(&start_device);
cudaEventCreate(&stop_device);
cudaEventRecord( start_device, 0 );

dim3 BS_MM(32,1,1);
dim3 GS_MM(1000,1,1);
MatrixMulXnor<<<GS_MM,BS_MM>>>(a_device_bin , b_device_bin , a_rows , MaxBS , b_cols ,
result_device , pitch_a_device_bin/sizeof(int) , pitch_b_device_bin/sizeof(int) ,
pitch_result_device/sizeof(float) , BINSIZE , a_cols);

cudaEventRecord( stop_device, 0 );
cudaEventSynchronize( stop_device );
cudaEventElapsedTime( &time_device, start_device, stop_device );
cudaEventDestroy( start_device );
cudaEventDestroy( stop_device );
cout<<"gputime="<<time_device<<"ms"<<endl;

cudaMemcpy2D(result,sizeof(float) *b_cols, result_device,pitch_result_device,sizeof(float) *b_cols , a_rows ,cudaMemcpyDeviceToDevice);

cudaFree(a_device);
cudaFree(b_device);
cudaFree(a_device_bin);
cudaFree(b_device_bin);
cudaFree(result_device);

//run in gpu warp in C code
MatrixMul_device(a_device,b_device,Matrixsize,Matrixsize,Matrixsize,result_device);

cudaMemcpy(result_host, result_device,sizeof(float) *Matrixsize * Matrixsize,cudaMemcpyDeviceToHost);
cudaFree(a_device);
cudaFree(b_device);
cudaFree(result_device);
// MatrixPrint(result_host,Matrixsize,Matrixsize);

//run in cpu
float *result_cpu;
result_cpu = (float*) malloc(sizeof(float) * Matrixsize * Matrixsize);
clock_t start_host = clock();
MatrixMul_host(a_host,Matrixsize,Matrixsize,b_host,Matrixsize,Matrixsize,result_cpu);
cout<<"cputime="<<(double)(clock() - start_host)/1000<<"ms"<<endl;
// MatrixPrint(result_cpu,Matrixsize,Matrixsize);

//compare value of gpu and cpu
float err=MatrixCompare(result_cpu,result_host,Matrixsize,Matrixsize);
cout<<"err in gpu and cpu = "<<err<<endl;

矩阵相乘（第“5×(1000的0次方)”讲）

__global__ void MatrixMulXnor(int *a,int *b,int a_rows,int a_cols,
int b_cols,float *result,int pitch_a,int pitch_b,
int pitch_result,int BINSIZE,int RealMidSize){
int tix=threadIdx.x;
// int tiy=threadIdx.y;
int bix=blockIdx.x;
// int biy=blockIdx.y;
int bdx=blockDim.x;
// int bdy=blockDim.y;
int gdx=gridDim.x;
// int gdy=gridDim.y;
extern __shared__ int a_shared[];
for(int j=tix;j<a_cols;j+=bdx){
// printf("i=%d ; j=%d\n",i,j);
a_shared[j]=a[bix*pitch_a+j];
}
__syncthreads();
int rest=(BINSIZE*a_cols-RealMidSize);

for(int j=tix;j<b_cols;j+=bdx){
// printf("i=%d ; j=%d\n",i,j);
int sum=0;
for(int k=0;k<a_cols;k++){
int bin=(a_shared[k]^b[k*pitch_b+j]);
int negnum=popcount(bin);
int posnum=BINSIZE-negnum;
//calculate ignores the rest of BINSIZE if the Matsize can't devided by BINSIZE ,it can cause err
//(10/00)'(01/00) should be 0000 but it is 0011,so 1+1 is trash in the result.and it mislead a_rows*b_cols times.
sum+=(posnum-negnum);
}
result[bix*pitch_result+j]=sum-rest;
}
}

矩阵相乘 倒数第二讲

int RectBlockSize = 16;
dim3 RectBlockNum_a_bin((a_rows-1)/RectBlockSize+1, (MaxBlocks-1)/RectBlockSize+1, 1);//with block multiply
dim3 RectBlockNum_b_bin((MaxBlocks-1)/RectBlockSize+1, (b_cols-1)/RectBlockSize+1, 1);
int *a_bin;
int *b_bin;
size_t Pitch_a_bin, Pitch_b_bin;
cudaMallocPitch((void**)&a_bin , &Pitch_a_bin , sizeof(int)*RectBlockSize*RectBlockNum_a_bin.y, RectBlockSize*RectBlockNum_a_bin.x);
cudaMallocPitch((void**)&b_bin , &Pitch_b_bin , sizeof(int)*RectBlockSize*RectBlockNum_b_bin.y, RectBlockSize*RectBlockNum_b_bin.x);
cudaMemset(a_bin, 0, Pitch_a_bin*RectBlockSize*RectBlockNum_a_bin.x);
cudaMemset(b_bin, 0, Pitch_b_bin*RectBlockSize*RectBlockNum_b_bin.x);
dim3 BS_BIN(512,1,1);
dim3 GS_BIN(6,1,1);
AMatrix2Bin<<< GS_BIN, BS_BIN >>>(a_copy, a_bin,
Pitch_a_copy/sizeof(float), Pitch_a_bin/sizeof(int), a_rows, MaxBlocks, BINSIZE);
BMatrix2Bin<<< GS_BIN, BS_BIN >>>(b_copy, b_bin,
Pitch_b_copy/sizeof(float), Pitch_b_bin/sizeof(int), b_cols, MaxBlocks, BINSIZE);
cudaFree(a_copy);
cudaFree(b_copy);

float *result_bin;//a_rows * b_cols
size_t Pitch_result_bin;
cudaMallocPitch((void**)&result_bin , &Pitch_result_bin , sizeof(float)*RectBlockSize*RectBlockNum_b_bin.y, RectBlockSize*RectBlockNum_a_bin.x);
const unsigned char __popcount_tab[] = {
0,1,1,2,1,2,2,3,1,2,2,3,2,3,3,4,1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,
1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,
1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,
2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,
1,2,2,3,2,3,3,4,2,3,3,4,3,4,4,5,2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,
2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,
2,3,3,4,3,4,4,5,3,4,4,5,4,5,5,6,3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,
3,4,4,5,4,5,5,6,4,5,5,6,5,6,6,7,4,5,5,6,5,6,6,7,5,6,6,7,6,7,7,8,
};
unsigned char *__popcount_tab_copy;
cudaMalloc((void**)&__popcount_tab_copy, sizeof(__popcount_tab));
cudaMemcpy(__popcount_tab_copy, __popcount_tab, sizeof(__popcount_tab), cudaMemcpyHostToDevice);

_global__ void MatrixMulXnor(int *a, int *b, float *result, unsigned char *__popcount_tab,
int pitch_a, int pitch_b, int pitch_result,
int midBlocks, int BINSIZE, int RealMidSize)

cudaEvent_t start_device, stop_device;
float time_device;
cudaEventCreate(&start_device);
cudaEventCreate(&stop_device);
cudaEventRecord(start_device, 0);
dim3 BS_MM(RectBlockSize, RectBlockSize, 1);
dim3 GS_MM(RectBlockNum_a_bin.x, RectBlockNum_b_bin.y, 1);
MatrixMulXnor<<< GS_MM, BS_MM >>>(a_bin, b_bin, result_bin, __popcount_tab_copy,
Pitch_a_bin/sizeof(int), Pitch_b_bin/sizeof(int), Pitch_result_bin/sizeof(float),
RectBlockNum_a_bin.y, BINSIZE, a_cols);
cudaEventRecord( stop_device, 0 );
cudaEventSynchronize( stop_device );
cudaEventElapsedTime( &time_device, start_device, stop_device );
cudaEventDestroy( start_device );
cudaEventDestroy( stop_device );
cout<<"gputime="<<time_device<<"ms"<<endl;
cudaMemcpy2D(result,sizeof(float) *b_cols, result_bin,Pitch_result_bin,sizeof(float) *b_cols , a_rows ,cudaMemcpyDeviceToDevice);
cudaFree(a_bin);
cudaFree(b_bin);
cudaFree(result_bin);

矩阵相乘 最后一讲

感谢大家的阅读，若有错误，欢迎批评职责，若有问题，欢迎交流。