#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda_runtime.h>
#define BLOCK_SIZE 16
static void HandleError(cudaError_t err, const char *file, int line)
{
if (err != cudaSuccess)
{
printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line);
exit(EXIT_FAILURE);
}
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
#define HANDLE_NULL( a ) {if ((a) == NULL) { \
printf("Host memory failed in %s at line %d\n", \
__FILE__, __LINE__); \
exit(EXIT_FAILURE); }}
static void GenerateNumbers(int *number, int size)
{
for (int i = 0; i < size; i++)
{
number[i] = rand() % 10;
}
}
static bool InitCUDA()
{
int count;
cudaGetDeviceCount(&count);
if (count == 0)
{
fprintf(stderr, "There is no device.\n");
return false;
}
int i;
for (i = 0; i < count; i++)
{
cudaDeviceProp prop;
if (cudaGetDeviceProperties(&prop, i) == cudaSuccess)
{
if (prop.major >= 1)
{
break;
}
}
}
if (i >= count)
{
fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
return false;
}
cudaSetDevice(i);
return true;
}
//多个block, 每一个block内256个线程,一个线程计算一个累加结果存放在global memory中
//每一个块都由块内threadIdx.x = 0的线程记录该块的启动时间和结束时间
__global__ static void sumOfSquares(int *num, int size, int* result, clock_t* time)
{
const int bid = blockIdx.x;
const int tid = threadIdx.x;
int sum = 0;
int index;
int totalThreadNum = gridDim.x * blockDim.x;
int gap = (size + totalThreadNum - 1) / totalThreadNum;
const int id = blockDim.x * blockIdx.x + tid;
if (tid == 0)
{
time[bid] = clock();
}
for (index = id * gap; index < (id + 1) * gap; index++)
{
if (index < size)
{
sum += num[index] * num[index];
}
}
result[id] = sum;
if (tid == 0)
{
time[bid + gridDim.x] = clock();
}
}
int main(int argc, char *argv[])
{
if (!InitCUDA())
{
return -1;
}
printf("CUDA initialized.\n");
//const int block_num = 32;
const int block_num = 16;
const int thread_num = 256;
const int DATA_SIZE = 1024 * 16 * 2;
int data[DATA_SIZE];
GenerateNumbers(data, DATA_SIZE);
int* gpudata, *result;
clock_t* devTime, gpuTime[2 * block_num], cpuTime;
HANDLE_ERROR(cudaMalloc((void **)&gpudata, sizeof(int)* DATA_SIZE));
HANDLE_ERROR(cudaMalloc((void**)&result, sizeof(int)* thread_num * block_num));
HANDLE_ERROR(cudaMalloc((void**)&devTime, sizeof(clock_t)* 2 * block_num));
HANDLE_ERROR(cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice));
sumOfSquares << <block_num, thread_num, 0 >> >(gpudata, DATA_SIZE, result, devTime);
long long sum = 0;
int gpuSum[thread_num * block_num];
HANDLE_ERROR(cudaMemcpy(gpuSum, result, sizeof(int)* block_num * thread_num, cudaMemcpyDeviceToHost));
HANDLE_ERROR(cudaMemcpy(gpuTime, devTime, sizeof(clock_t)* 2 * block_num, cudaMemcpyDeviceToHost));
cudaFree(gpudata);
cudaFree(result);
cudaFree(devTime);
for (int i = 0; i < block_num * thread_num; i++)
{
sum += gpuSum[i];
}
clock_t minStart = gpuTime[0];
clock_t maxEnd = gpuTime[block_num];
for (int i = 1; i < block_num; i++)
{
if (minStart > gpuTime[i])
{
minStart = gpuTime[i];
}
if (maxEnd < gpuTime[i + block_num])
{
maxEnd = gpuTime[i + block_num];
}
}
printf("sum (GPU): %ld, time: %d\n", sum, maxEnd - minStart);
long long sumCPU = 0;
cpuTime = clock();
for (int i = 0; i < DATA_SIZE; i++)
{
sumCPU += data[i] * data[i];
}
cpuTime = clock() - cpuTime;
printf("sum (CPU): %ld, time: %d\n", sumCPU, cpuTime);
printf("Result %s\n", sum == sumCPU ? "OK" : "Wrong");
//remember to release the device
cudaDeviceReset();
return 0;
}
cuda编程入门示例13
最新推荐文章于 2023-12-30 23:22:44 发布