#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <chrono>
#include <thread>
#define MATRIX_SIZE 1024
#define BLOCK_SIZE 16
__global__ void matrixMultiply(float* A, float* B, float* C, int size) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < size && col < size) {
float sum = 0.0f;
for (int k = 0; k < size; k++) {
sum += A[row * size + k] * B[k * size + col];
}
C[row * size + col] = sum;
}
}
__global__ void vectorAdd(float* A, float* B, float* C, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
C[idx] = A[idx] + B[idx];
}
}
__global__ void reduceSum(float* input, float* output, int size) {
extern __shared__ float sdata[];
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
sdata[tid] = (i < size) ? input[i] : 0.0f;
__syncthreads();
for (unsigned int s = blockDim.x / 2; s >
0; s >>= 1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
if (tid == 0) {
output[blockIdx.x] = sdata[0];
}
}
#define CHECK_CUDA_ERROR(err) \
if (err != cudaSuccess) { \
printf("CUDA Error: %s\n", cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
}
void printGPUInfo() {
int deviceCount;
cudaGetDeviceCount(&deviceCount);
printf("Number of CUDA devices: %d\n", deviceCount);
for (int i = 0; i < deviceCount; i++) {
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, i);
printf("\nDevice %d: %s\n", i, prop.name);
printf(" Compute capability: %d.%d\n", prop.major, prop.minor);
printf(" Total global memory: %.2f GB\n", prop.totalGlobalMem / 1024.0 / 1024.0 / 1024.0);
printf(" Multiprocessors: %d\n", prop.multiProcessorCount);
printf(" Max threads per block: %d\n", prop.maxThreadsPerBlock);
printf(" Max threads per multiprocessor: %d\n", prop.maxThreadsPerMultiProcessor);
}
printf("\n");
}
int main() {
printf("Starting CUDA Stress Test for 2 minutes...\n");
printGPUInfo();
int device = 0;
cudaError_t err = cudaSetDevice(device);
CHECK_CUDA_ERROR(err);
size_t matrixSize = MATRIX_SIZE * MATRIX_SIZE * sizeof(float);
size_t vectorSize = MATRIX_SIZE * MATRIX_SIZE * sizeof(float);
float *h_A = (float*)malloc(matrixSize);
float *h_B = (float*)malloc(matrixSize);
float *h_C = (float*)malloc(matrixSize);
for (int i = 0; i < MATRIX_SIZE * MATRIX_SIZE; i++) {
h_A[i] = rand() / (float)RAND_MAX;
h_B[i] = rand() / (float)RAND_MAX;
}
float *d_A, *d_B, *d_C;
err = cudaMalloc(&d_A, matrixSize);
CHECK_CUDA_ERROR(err);
err = cudaMalloc(&d_B, matrixSize);
CHECK_CUDA_ERROR(err);
err = cudaMalloc(&d_C, matrixSize);
CHECK_CUDA_ERROR(err);
err = cudaMemcpy(d_A, h_A, matrixSize, cudaMemcpyHostToDevice);
CHECK_CUDA_ERROR(err);
err = cudaMemcpy(d_B, h_B, matrixSize, cudaMemcpyHostToDevice);
CHECK_CUDA_ERROR(err);
dim3 threadsPerBlock(BLOCK_SIZE, BLOCK_SIZE);
dim3 blocksPerGrid((MATRIX_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE,
(MATRIX_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE);
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
auto startTime = std::chrono::high_resolution_clock::now();
auto endTime = startTime + std::chrono::minutes(2);
int iteration = 0;
float totalGFlops = 0.0f;
printf("Running stress test for 2 minutes...\n");
printf("Iteration | Time (ms) | GFLOP/s\n");
printf("----------|-----------|---------\n");
while (std::chrono::high_resolution_clock::now() < endTime) {
iteration++;
cudaEventRecord(start);
matrixMultiply<<
<blocksPerGrid, threadsPerBlock>>
>
(d_A, d_B, d_C, MATRIX_SIZE);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
float flops = 2.0f * MATRIX_SIZE * MATRIX_SIZE * MATRIX_SIZE;
float gflops = (flops / milliseconds) / 1e6;
totalGFlops += gflops;
if (iteration % 10 == 0) {
printf("%9d | %9.3f | %7.2f\n", iteration, milliseconds, gflops);
}
if (iteration % 50 == 0) {
err = cudaDeviceSynchronize();
CHECK_CUDA_ERROR(err);
for (int i = 0; i < MATRIX_SIZE * MATRIX_SIZE; i++) {
h_A[i] = rand() / (float)RAND_MAX;
h_B[i] = rand() / (float)RAND_MAX;
}
cudaMemcpy(d_A, h_A, matrixSize, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, matrixSize, cudaMemcpyHostToDevice);
}
}
float avgGFlops = totalGFlops / iteration;
auto totalDuration = std::chrono::duration_cast<std::chrono::milliseconds>
(
std::chrono::high_resolution_clock::now() - startTime);
printf("\nTest completed successfully!\n");
printf("Total iterations: %d\n", iteration);
printf("Total time: %.2f seconds\n", totalDuration.count() / 1000.0);
printf("Average performance: %.2f GFLOP/s\n", avgGFlops);
free(h_A);
free(h_B);
free(h_C);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaDeviceReset();
printf("CUDA stress test completed successfully!\n");
return 0;
}
学习助手
订阅
