运行耗时比较
第一次(完成功能)
运行结果
完整代码
/* 利用CUDA计算1M大小数据的立方和 */
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda_runtime.h>
//1M
#define DATA_SIZE 1048576
#define THREAD_NUM 256
// 存储使用到的数据
int data[DATA_SIZE] = { 0 };
// 计算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
//表示目前的 thread 是第几个 thread(由 0 开始计算)
const int tid = threadIdx.x;
//计算每个线程需要完成的量
const int size = DATA_SIZE / THREAD_NUM;
int sum = 0;
// 各线程计算属于自己的部分
for (int i = tid * size; i < (tid + 1) * size; i++)
{
sum += num[i] * num[i] * num[i];
}
// 将各线程"立方和"存储
result[tid] = sum;
}
int main()
{
// 设置显卡(若无显卡则程序会失败,此为演示则不进行判断)
cudaSetDevice(0);
// 对需要使用的数据填充随机数
for (int i = 0; i < DATA_SIZE; i++) { data[i] = rand() % 10; }
// 把数据复制到显卡内存中
int* gpudata, *result;
cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM);
cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
// 使用cuda事件来计算耗时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
// 在CUDA 中执行函数 语法:函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
sumOfSquares << < 1, THREAD_NUM, 0 >> > (gpudata, result);
// 同步及记录事件
cudaEventRecord(stop);
cudaEventSynchronize(stop);
// 把结果从显示芯片复制回主内存
int sum[THREAD_NUM] = { 0 };
cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM, cudaMemcpyDeviceToHost);
//Free
if (gpudata) { cudaFree(gpudata); }gpudata = nullptr;
if (result) { cudaFree(result); }result = nullptr;
// 计算事件之间经过的时间
float time = 0;
cudaEventElapsedTime(&time, start, stop);
int final_sum = 0;
for (int i = 0; i < THREAD_NUM; i++) { final_sum += sum[i]; }
printf("GPUsum: %d , time: %gms\n\n", final_sum, time);
final_sum = 0;
// cpu计算结果核对
for (int i = 0; i < DATA_SIZE; i++) { final_sum += data[i] * data[i] * data[i]; }
printf("CPUsum: %d \n", final_sum);
return 0;
}
第二次(使用"内存连续"方法来优化核函数)
运行结果
完整代码(在第一次代码中仅修改了核函数,基本是for循环修改)
// 计算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
//表示目前的 thread 是第几个 thread(由 0 开始计算)
const int tid = threadIdx.x;
//计算每个线程需要完成的量
int sum = 0;
// 各线程计算属于自己的部分
for (int i = tid ; i < DATA_SIZE; i+= THREAD_NUM)
{
sum += num[i] * num[i] * num[i];
}
// 将各线程"立方和"存储
result[tid] = sum;
}
第三次(使用"线程块"方法来优化核函数)
运行结果
完整代码
/* 利用CUDA计算1M大小数据的立方和 */
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda_runtime.h>
//1M
#define DATA_SIZE 1048576
#define THREAD_NUM 256
#define BLOCK_NUM 32
// 存储使用到的数据
int data[DATA_SIZE] = { 0 };
// 计算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
//表示目前的 thread 是第几个 thread(由 0 开始计算)
const int tid = threadIdx.x;
//表示目前的 thread 属于第几个 block(由 0 开始计算)
const int bid = blockIdx.x;
//计算每个线程需要完成的量
int sum = 0;
// 各线程计算属于自己的部分
for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM)
{
sum += num[i] * num[i] * num[i];
}
// 将各线程"立方和"存储
result[bid * THREAD_NUM + tid] = sum;
}
int main()
{
// 设置显卡(若无显卡则程序会失败,此为演示则不进行判断)
cudaSetDevice(0);
// 对需要使用的数据填充随机数
for (int i = 0; i < DATA_SIZE; i++) { data[i] = rand() % 10; }
// 把数据复制到显卡内存中
int* gpudata, *result;
cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
cudaMalloc((void**)&result, sizeof(int)*BLOCK_NUM*THREAD_NUM);
cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
// 使用cuda事件来计算耗时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
// 在CUDA 中执行函数 语法:函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
sumOfSquares << < BLOCK_NUM, THREAD_NUM, 0 >> > (gpudata, result);
// 同步及记录事件
cudaEventRecord(stop);
cudaEventSynchronize(stop);
// 把结果从显示芯片复制回主内存
int sum[BLOCK_NUM*THREAD_NUM] = { 0 };
cudaMemcpy(&sum, result, sizeof(int)* BLOCK_NUM*THREAD_NUM, cudaMemcpyDeviceToHost);
//Free
if (gpudata) { cudaFree(gpudata); }gpudata = nullptr;
if (result) { cudaFree(result); }result = nullptr;
// 计算事件之间经过的时间
float time = 0;
cudaEventElapsedTime(&time, start, stop);
int final_sum = 0;
for (int i = 0; i < BLOCK_NUM * THREAD_NUM; i++) { final_sum += sum[i]; }
printf("GPUsum: %d , time: %gms\n\n", final_sum, time);
final_sum = 0;
// cpu计算结果核对
for (int i = 0; i < DATA_SIZE; i++) { final_sum += data[i] * data[i] * data[i]; }
printf("CPUsum: %d \n", final_sum);
return 0;
}
第四次(使用"共享内存"方法来优化核函数)
运行结果
完整代码
/* 利用CUDA计算1M大小数据的立方和 */
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda_runtime.h>
//1M
#define DATA_SIZE 1048576
#define THREAD_NUM 256
#define BLOCK_NUM 32
// 存储使用到的数据
int data[DATA_SIZE] = { 0 };
// 计算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
//声明一块共享内存
extern __shared__ int shared[];
//表示目前的 thread 是第几个 thread(由 0 开始计算)
const int tid = threadIdx.x;
shared[tid] = 0;
//表示目前的 thread 属于第几个 block(由 0 开始计算)
const int bid = blockIdx.x;
// 各线程计算属于自己的部分
for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM)
{
shared[tid] += num[i] * num[i] * num[i];
}
//同步 保证每个 thread 都已经把结果写到 shared[tid] 里面
__syncthreads();
//使用线程0完成加和
if (tid == 0)
{
for (int i = 1; i < THREAD_NUM; ++i)
{
shared[0] += shared[i];
}
result[bid] = shared[0];
}
}
int main()
{
// 设置显卡(若无显卡则程序会失败,此为演示则不进行判断)
cudaSetDevice(0);
// 对需要使用的数据填充随机数
for (int i = 0; i < DATA_SIZE; i++) { data[i] = rand() % 10; }
// 把数据复制到显卡内存中
int* gpudata, *result;
cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
cudaMalloc((void**)&result, sizeof(int)*BLOCK_NUM);
cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
// 使用cuda事件来计算耗时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
// 在CUDA 中执行函数 语法:函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
sumOfSquares << < BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int) >> > (gpudata, result);
// 同步及记录事件
cudaEventRecord(stop);
cudaEventSynchronize(stop);
// 把结果从显示芯片复制回主内存
int sum[BLOCK_NUM] = { 0 };
cudaMemcpy(&sum, result, sizeof(int)* BLOCK_NUM, cudaMemcpyDeviceToHost);
//Free
if (gpudata) { cudaFree(gpudata); }gpudata = nullptr;
if (result) { cudaFree(result); }result = nullptr;
// 计算事件之间经过的时间
float time = 0;
cudaEventElapsedTime(&time, start, stop);
int final_sum = 0;
for (int i = 0; i < BLOCK_NUM; i++) { final_sum += sum[i]; }
printf("GPUsum: %d , time: %gms\n\n", final_sum, time);
final_sum = 0;
// cpu计算结果核对
for (int i = 0; i < DATA_SIZE; i++) { final_sum += data[i] * data[i] * data[i]; }
printf("CPUsum: %d \n", final_sum);
return 0;
}
关注
笔者 - 东旭