核函数外部并行(可分为5类):
1. 核函数计算与数据传输之间的并行
2. 主机计算与数据传输之间的并行
3. 不同的数据传输之间的并行
4. 核函数计算与主机计算之间的并行
5. 不同核函数之间的并行
以下代码示例为第5类情况分别对RTX3080Ti在同时并行1-128个核函数的情况进行测试:
kernel-kernel_1024.cu
#include "error.cuh"
#include <math.h>
#include <stdio.h>
#ifdef USE_DP
typedef double real;
#else
typedef float real;
#endif
const int NUM_REPEATS = 10;
const int N1 = 1024; //每个核函数使用1024个线程,后改为2048进行第二次测试
const int MAX_NUM_STREAMS = 128; //分别递增测试到最大同时并行128个kernel函数
const int N = N1 * MAX_NUM_STREAMS;
const int M = sizeof(real) * N;
const int block_size = 128;
const int grid_size = (N1 - 1) / block_size + 1;
cudaStream_t streams[MAX_NUM_STREAMS];
void timing(const real *d_x, const real *d_y, real *d_z, const int num);
int main(void)
{
real *h_x = (real*) malloc(M);
real *h_y = (real*) malloc(M);
for (int n = 0; n < N; ++n)
{
h_x[n] = 1.23;
h_y[n] = 2.34;
}
real *d_x, *d_y, *d_z;
CHECK(cudaMalloc(&d_x, M));
CHECK(cudaMalloc(&d_y, M));
CHECK(cudaMalloc(&d_z, M));
CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
for (int n = 0 ; n < MAX_NUM_STREAMS; ++n)
{
CHECK(cudaStreamCreate(&(streams[n])));
}
for (int num = 1; num <= MAX_NUM_STREAMS; ++num)
{
timing(d_x, d_y, d_z, num); //num是流编号,最大128
}
for (int n = 0 ; n < MAX_NUM_STREAMS; ++n)
{
CHECK(cudaStreamDestroy(streams[n]));
}
free(h_x);
free(h_y);
CHECK(cudaFree(d_x));
CHECK(cudaFree(d_y));
CHECK(cudaFree(d_z));
return 0;
}
void __global__ add(const real *d_x, const real *d_y, real *d_z)
{
const int n = blockDim.x * blockIdx.x + threadIdx.x;
if (n < N1) //每个核函数使用N1个线程
{
for (int i = 0; i < 100000; ++i) //注意,这个for循环是故意让每个线程重复100000加法,拖延时间
{
d_z[n] = d_x[n] + d_y[n];
}
}
}
void timing(const real *d_x, const real *d_y, real *d_z, const int num) //num是流编号
{
float t_sum = 0;
float t2_sum = 0;
for (int repeat = 0; repeat <= NUM_REPEATS; ++repeat)
{
cudaEvent_t start, stop;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventRecord(start));
cudaEventQuery(start);
for (int n = 0; n < num; ++n) //num是流编号
{
int offset = n * N1; //每个核函数使用N1个线程
add<<<grid_size, block_size, 0, streams[n]>>>
(d_x + offset, d_y + offset, d_z + offset);
}
CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
float elapsed_time;
CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
if (repeat > 0)
{
t_sum += elapsed_time;
t2_sum += elapsed_time * elapsed_time;
}
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
}
const float t_ave = t_sum / NUM_REPEATS;
const float t_err = sqrt(t2_sum / NUM_REPEATS - t_ave * t_ave);
printf("%g\n", t_ave);
}
头文件error.cuh为错误检查宏CHECK函数:
#pragma once
#include <stdio.h>
#define CHECK(call) \
do \
{ \
const cudaError_t error_code = call; \
if (error_code != cudaSuccess) \
{ \
printf("CUDA Error:\n"); \
printf(" File: %s\n", __FILE__); \
printf(" Line: %d\n", __LINE__); \
printf(" Error code: %d\n", error_code); \
printf(" Error text: %s\n", \
cudaGetErrorString(error_code)); \
exit(1); \
} \
} while (0)
编译并运行:
nvcc kernel-kernel_1024.cu -o kernel-kernel_1024
./kernel-kernel_1024 > k-k_1024.dat
nvcc kernel-kernel_2048.cu -o kernel-kernel_2048
./kernel-kernel_2048 > k-k_2048.dat
分别对 k-k_1024.dat 和 k-k_2048.dat 两个结果文件进行绘图,结果如下:
可以看出,我使用的 RTX3080Ti 在使用32个SM(流多处理器)以下,每个SM常驻1024个线程以下时,计算时间不变。(因为在横坐标CUDA stream低于32时,每个SM驻1024个线程时间均为4ms;在横坐标CUDA stream低于16时,每个SM驻2048个线程时间均为4ms)
然而我的 RTX3080Ti 理论结果应该是有80个SM(流多处理器),每个SM理论上可以常驻1536个线程。
$ ./query
Device id: 0
Device name: NVIDIA GeForce RTX 3080 Ti
Compute capability: 8.6
Amount of global memory: 11.7689 GB
Amount of constant memory: 64 KB
Maximum grid size: 2147483647 65535 65535
Maximum block size: 1024 1024 64
Number of SMs: 80
Maximum amount of shared memory per block: 48 KB
Maximum amount of shared memory per SM: 100 KB
Maximum number of registers per block: 64 K
Maximum number of registers per SM: 64 K
Maximum number of threads per block: 1024
Maximum number of threads per SM: 1536
(base)