#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <time.h>
#include <stdlib.h> #define MAX 120
#define MIN 0 cudaError_t addWithCudaStream(int *c, const int *a, const int *b, size_t size,
float* etime);
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size,
float* etime, int type);
__global__ void addKernel(int *c, const int *a, const int *b) {
int i = blockIdx.x;
c[i] = a[i] + b[i];
} __global__ void addKernelThread(int *c, const int *a, const int *b) {
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
int main() {
const int arraySize = ;
srand((unsigned) time(NULL));
int a[arraySize] = { , , , , };
int b[arraySize] = { , , , , }; for (int i = ; i < arraySize; i++) {
a[i] = rand() % (MAX + - MIN) + MIN;
b[i] = rand() % (MAX + - MIN) + MIN;
}
int c[arraySize] = { };
// Add vectors in parallel.
cudaError_t cudaStatus;
int num = ;
cudaDeviceProp prop;
cudaStatus = cudaGetDeviceCount(&num);
for (int i = ; i < num; i++) {
cudaGetDeviceProperties(&prop, i);
} float time;
cudaStatus = addWithCudaStream(c, a, b, arraySize, &time);
printf("Elasped time of stream is : %f \n", time);
printf("{%d,%d,%d,%d,%d} + {%d,%d,%d,%d,%d} = {%d,%d,%d,%d,%d}\n",
a[arraySize - - ], a[arraySize - - ], a[arraySize - - ],
a[arraySize - - ], a[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], b[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], c[arraySize - - ], c[arraySize - - ],
c[arraySize - - ], c[arraySize - - ], c[arraySize - - ]);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCudaStream failed!");
return ;
}
cudaStatus = addWithCuda(c, a, b, arraySize, &time, );
printf("Elasped time of Block is : %f \n", time);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCudaStream failed!");
return ;
}
printf("{%d,%d,%d,%d,%d} + {%d,%d,%d,%d,%d} = {%d,%d,%d,%d,%d}\n",
a[arraySize - - ], a[arraySize - - ], a[arraySize - - ],
a[arraySize - - ], a[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], b[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], c[arraySize - - ], c[arraySize - - ],
c[arraySize - - ], c[arraySize - - ], c[arraySize - - ]); cudaStatus = addWithCuda(c, a, b, arraySize, &time, );
printf("Elasped time of thread is : %f \n", time);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCudaStream failed!");
return ;
}
printf("{%d,%d,%d,%d,%d} + {%d,%d,%d,%d,%d} = {%d,%d,%d,%d,%d}\n",
a[arraySize - - ], a[arraySize - - ], a[arraySize - - ],
a[arraySize - - ], a[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], b[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], c[arraySize - - ], c[arraySize - - ],
c[arraySize - - ], c[arraySize - - ], c[arraySize - - ]); cudaStatus = addWithCudaStream(c, a, b, arraySize, &time);
printf("Elasped time of stream is : %f \n", time);
printf("{%d,%d,%d,%d,%d} + {%d,%d,%d,%d,%d} = {%d,%d,%d,%d,%d}\n",
a[arraySize - - ], a[arraySize - - ], a[arraySize - - ],
a[arraySize - - ], a[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], b[arraySize - - ], b[arraySize - - ],
b[arraySize - - ], c[arraySize - - ], c[arraySize - - ],
c[arraySize - - ], c[arraySize - - ], c[arraySize - - ]);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCudaStream failed!");
return ;
}
// cudaThreadExit must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaThreadExit();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaThreadExit failed!");
return ;
}
return ;
}
// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCudaStream(int *c, const int *a, const int *b, size_t size,
float* etime) {
int *dev_a = ;
int *dev_b = ;
int *dev_c = ;
clock_t start, stop;
float time;
cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice();
if (cudaStatus != cudaSuccess) {
fprintf(stderr,
"cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**) &dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**) &dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**) &dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int),
cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int),
cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStream_t stream[];
for (int i = ; i < ; i++) {
cudaStreamCreate(&stream[i]); //创建流
}
// Launch a kernel on the GPU with one thread for each element.
for (int i = ; i < ; i++) {
addKernel<<<, , , stream[i]>>>(dev_c + i, dev_a + i, dev_b + i); //执行流
}
start = clock();
cudaDeviceSynchronize();
stop = clock();
time = (float) (stop - start) / CLOCKS_PER_SEC;
*etime = time;
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr,
"cudaThreadSynchronize returned error code %d after launching addKernel!\n",
cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int),
cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error: for (int i = ; i < ; i++) {
cudaStreamDestroy(stream[i]); //销毁流
}
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size,
float * etime, int type) {
int *dev_a = ;
int *dev_b = ;
int *dev_c = ;
clock_t start, stop;
float time;
cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice();
if (cudaStatus != cudaSuccess) {
fprintf(stderr,
"cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}
// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**) &dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**) &dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
cudaStatus = cudaMalloc((void**) &dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}
// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int),
cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int),
cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
} if (type == ) {
start = clock();
addKernel<<<size, >>>(dev_c, dev_a, dev_b);
} else {
start = clock();
addKernelThread<<<, size>>>(dev_c, dev_a, dev_b);
}
stop = clock();
time = (float) (stop - start) / CLOCKS_PER_SEC;
*etime = time;
// cudaThreadSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaThreadSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr,
"cudaThreadSynchronize returned error code %d after launching addKernel!\n",
cudaStatus);
goto Error;
}
// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int),
cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}
Error: cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);
return cudaStatus;
}
如上文的实现程序，使用了thread并行，block并行，stream并行三种，使用三种方法法进行了五次计算，发现stream第一次计算时会出错，调用的子程序没有变化，没有搞懂？
Elasped time of stream is : 0.000006
{47,86,67,35,16} + {114,39,110,20,101} = {158,123,92,107,127}
Elasped time of Block is : 0.000006
{47,86,67,35,16} + {114,39,110,20,101} = {161,125,177,55,117}
Elasped time of stream is : 0.000008
{47,86,67,35,16} + {114,39,110,20,101} = {161,125,177,55,117}
Elasped time of thread is : 0.000004
{47,86,67,35,16} + {114,39,110,20,101} = {161,125,177,55,117}
Elasped time of stream is : 0.000007
{47,86,67,35,16} + {114,39,110,20,101} = {161,125,177,55,117}