在2.0中进行了用一维网格和块对一维向量进行了求和。
在2.1中例化了二维的网格和块。
接下来进行2维网络（grid）和2维块（block）对矩阵进行求和。
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <windows.h>typedef unsigned long DWORD;#define CHECK(call) \{\const cudaError_t error = call; \if (error != cudaSuccess)\{\printf("Error: %s: %d\n", __FILE__, __LINE__);\printf("code :%d reason :%s\n", error , cudaGetErrorString(error));\exit(1);\}\
}void checkResult(float *hostRef, float *gpuRef, const int N)
{double epsilon = 1.0E-8;bool match = 1;for (int i = 0; i < N; i++){if (abs(hostRef[i] - gpuRef[i])> epsilon){match = 0;printf("Array do not match\n");printf("host %5.2f gpu % 5.2f at current %d\n", hostRef[i], gpuRef[i], i);break;}}if (match) printf("array matches\n");
}void initialData(float *ip, int size)
{time_t t;srand((unsigned int) time(&t));for (int i = 0; i < size; i++) {ip[i] = (float) (rand() & 0xff) / 10.0f;}
}void sumMatrixOnHost(float *A, float *B, float *C, const int nx, const int ny){float *ia = A;float *ib = B;float *ic = C;for (int iy = 0; iy < ny; iy++){for (int ix =0; ix < nx; ix++){ic[ix] = ia[ix] + ib[ix];}ia += nx; ib += nx;ic += nx;}
}__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx, int ny){unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;unsigned int idx = iy*nx + ix;if (ix < nx && iy < ny){MatC[idx] = MatA[idx] + MatB[idx];}
}int main(int argc , char **argv)
{printf("%s starting\n", argv[0]);int dev = 0;cudaDeviceProp deviceprop;CHECK(cudaGetDeviceProperties(&deviceprop,dev));printf("Using Device %d : %s\n", dev, deviceprop.name);CHECK(cudaSetDevice(dev));//set up dataint nx = 1<<14;int ny = 1<<14;int nxy = nx * ny;size_t nBytes = nxy * sizeof(float);printf("matrix size %d %d\n", nx, ny);float *h_A, *h_B, *hostRef, *gpuRef;h_A = (float *) malloc (nBytes);h_B = (float *) malloc (nBytes);hostRef = (float *) malloc (nBytes);gpuRef = (float *) malloc (nBytes);initialData(h_A, nxy);initialData(h_B, nxy);memset(hostRef,0, nBytes);memset(gpuRef,0, nBytes);// malloc device global memoryfloat *d_MatA, *d_MatB, *d_MatC;cudaMalloc((float**)&d_MatA, nBytes);cudaMalloc((float**)&d_MatB, nBytes);cudaMalloc((float**)&d_MatC, nBytes);//transfer data from host to devicecudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice);int dimx = 32;int dimy = 32;dim3 block(dimx, dimy);dim3 grid((nx + block.x - 1)/block.x, (ny + block.y - 1)/block.y);cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start);sumMatrixOnGPU2D<<<grid,block>>>(d_MatA, d_MatB, d_MatC, nx, ny);cudaDeviceSynchronize();cudaEventRecord(stop);cudaEventSynchronize(stop);float milliseconds = 0;cudaEventElapsedTime(&milliseconds, start, stop);printf("execution config <<<(%d,%d), (%d,%d)>>>\n", grid.x,grid.y, block.x, block.y);printf("Kernel execution time: %f ms\n", milliseconds);cudaEventDestroy(start);cudaEventDestroy(stop);//copy kernel result back to hostcudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost);sumMatrixOnHost(h_A, h_B, hostRef, nx,ny);checkResult(hostRef, gpuRef, nxy);cudaFree(d_MatA);cudaFree(d_MatB);cudaFree(d_MatC);free(h_A);free(h_B);free(hostRef);free(gpuRef);return 0;
}
基本流程和1维向量求和类似
输出结果：
Using Device 0 : NVIDIA GeForce RTX 4090
matrix size 16384 16384
execution config <<<(512,512), (32,32)>>>
Kernel execution time: 5.351136 ms
array matches
block的尺寸为32x32。//block(dimx,dimy)定义的。
改变block尺寸为32x16:
execution config <<<(512,1024), (32,16)>>>
Kernel execution time: 3.778752 ms
进一步改变block尺寸为16x16:
execution config <<<(1024,1024), (16,16)>>>
Kernel execution time: 3.712736 ms
在之前尝试使用nvprof测试kernl性能时，report
======= Warning: nvprof is not supported on devices with compute capability 8.0 and higher.
参考 https://blog.csdn.net/TH_NUM/article/details/109952643 使用nsys
将C:\Program Files\NVIDIA Corporation\Nsight Systems 2024.5.1\target-windows-x64加入环境变量即可
nsys profile --stats=true .\sum_matrix_on_gpu_timer.exe
输出：
Collecting data...
Generating 'C:\Users\ADMINI~1\AppData\Local\Temp\nsys-report-ffa3.qdstrm'
[1/8] [========================100%] report2.nsys-rep
[2/8] [========================100%] report2.sqlite
[3/8] Executing 'nvtx_sum' stats report
SKIPPED: C:\Users\Administrator\Desktop\edward_temp\chapter2\report2.sqlite does not contain NV Tools Extension (NVTX) data.
[4/8] Executing 'osrt_sum' stats report
SKIPPED: No data available.
[5/8] Executing 'cuda_api_sum' stats reportTime (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name-------- --------------- --------- ----------- ---------- -------- --------- ----------- ----------------------93.3 321764988 3 107254996.0 91069908.0 83897570 146797510 34432084.1 cudaMemcpy4.0 13772507 3 4590835.7 4393180.0 3984976 5394351 725179.5 cudaFree1.5 5118078 3 1706026.0 1249576.0 819401 3049101 1182856.9 cudaMalloc1.0 3496955 1 3496955.0 3496955.0 3496955 3496955 0.0 cudaDeviceSynchronize0.1 459711 1 459711.0 459711.0 459711 459711 0.0 cudaLaunchKernel0.0 49593 2 24796.5 24796.5 707 48886 34067.7 cudaEventCreate0.0 22341 1 22341.0 22341.0 22341 22341 0.0 cuLibraryUnload0.0 18196 2 9098.0 9098.0 7920 10276 1665.9 cudaEventRecord0.0 15060 1 15060.0 15060.0 15060 15060 0.0 cudaEventSynchronize0.0 1961 1 1961.0 1961.0 1961 1961 0.0 cuCtxSynchronize0.0 1434 1 1434.0 1434.0 1434 1434 0.0 cuModuleGetLoadingMode0.0 1012 2 506.0 506.0 205 807 425.7 cudaEventDestroy 0.0 181 1 181.0 181.0 181 181 0.0 cuDeviceGetLuid[6/8] Executing 'cuda_gpu_kern_sum' stats reportTime (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name-------- --------------- --------- --------- --------- -------- -------- ----------- -----------------------------------------------------100.0 3453326 1 3453326.0 3453326.0 3453326 3453326 0.0 sumMatrixOnGPU2D(float *, float *, float *, int, int)[7/8] Executing 'cuda_gpu_mem_time_sum' stats reportTime (%) Total Time (ns) Count Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Operation-------- --------------- ----- ---------- ---------- -------- -------- ----------- ----------------------------68.3 180949528 2 90474764.0 90474764.0 89939258 91010270 757319.8 [CUDA memcpy Host-to-Device]31.7 83834368 1 83834368.0 83834368.0 83834368 83834368 0.0 [CUDA memcpy Device-to-Host][8/8] Executing 'cuda_gpu_mem_size_sum' stats reportTotal (MB) Count Avg (MB) Med (MB) Min (MB) Max (MB) StdDev (MB) Operation---------- ----- -------- -------- -------- -------- ----------- ----------------------------2147.484 2 1073.742 1073.742 1073.742 1073.742 0.000 [CUDA memcpy Host-to-Device]1073.742 1 1073.742 1073.742 1073.742 1073.742 0.000 [CUDA memcpy Device-to-Host]Generated:C:\Users\Administrator\Desktop\edward_temp\chapter2\report2.nsys-repC:\Users\Administrator\Desktop\edward_temp\chapter2\report2.sqlite
