说明
样例的内容是每个CPU启用一个线程,执行一个核函数,然后这些线程平均分配给n个GPU上面运行,采用多线程技术实现。
代码
#include <cuda_runtime.h>
#include <pthread.h> //多线程
#include <unistd.h>
#include <stdio.h> // stdio functions are used since C++ streams aren't necessarily thread safe
#include <pthread.h>
#define address unsigned long long int //实现指针与整数的灵活转变
int* finish_code; //检查单独的一个线程是否执行完毕,为0代表执行完毕
// 获取CPU数量 Linux操作系统使用
int Get_CPU_Number_Linux() {
// #include <unistd.h>
return sysconf( _SC_NPROCESSORS_CONF);
}
// 获取GPU数量 不限操作系统
int Get_GPU_Number(){
// #include <cuda_runtime.h>
int num = 0;
cudaGetDeviceCount(&num);
return num;
}
//核函数,执行数组每个元素的固定值增加
__global__ void kernelAddConstant(int *g_a, int b) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
g_a[idx] += b;
}
// 线程运行函数
void* exec(void* args){
//进行函数声明,否则跨线程无法调用GPU核函数
__global__ void kernelAddConstant(int *g_a, int b);
// 将void*指针强制转为address*指针,存储的内容为若干整数,这些整数的数值即为地址
//addr[0] = cpu_thread_id; addr[1] = n
//addr[2] = nbytes; addr[3] = b; addr[4] = data
address* addr = (address*)args;
// 传入的整数强制转换为指针的地址
int cpu_thread_id = ((int*)(addr[0]))[0]; //获得CPU线程索引
int num_cpu_threads = Get_CPU_Number_Linux(); //获得CPU线程数量
int n = ((int*)(addr[1]))[0]; //获得GPU线程总数
int nbytes = ((int*)(addr[2]))[0]; //获得数组占用内存大小
int b = ((int*)(addr[3]))[0]; //获得数组元素自增值
int* data = ((int*)(addr[4])); //获得数组数据
// set and check the CUDA device for this CPU thread
int gpu_id = -1;
cudaSetDevice(cpu_thread_id %Get_GPU_Number()); // "% num_gpus" allows more CPU threads than GPU devices
cudaGetDevice(&gpu_id);
printf("CPU thread %d (of %d) uses CUDA device %d\n", cpu_thread_id,num_cpu_threads, gpu_id);
int *d_a = 0; // pointer to memory on the device associated with this CPU thread
int *sub_a =
data +
cpu_thread_id * n /
num_cpu_threads; // pointer to this CPU thread's portion of data
unsigned int nbytes_per_kernel = nbytes / num_cpu_threads;
dim3 gpu_threads(128); // 128 threads per block
dim3 gpu_blocks(n / (gpu_threads.x * num_cpu_threads));
cudaMalloc((void **)&d_a, nbytes_per_kernel);
cudaMemset(d_a, 0, nbytes_per_kernel);
cudaMemcpy(d_a, sub_a, nbytes_per_kernel, cudaMemcpyHostToDevice);
kernelAddConstant<< <gpu_blocks, gpu_threads>> >(d_a, b);
cudaMemcpy(sub_a, d_a, nbytes_per_kernel, cudaMemcpyDeviceToHost);
cudaFree(d_a);
finish_code[cpu_thread_id] = 0;
printf("thread %d finish!\n", cpu_thread_id);
}
//检查数组的每个值是不是index + b
int correctResult(int *data, const int n, const int b) {
for (int i = 0; i < n; i++)
if (data[i] != i + b) return 0;
return 1;
}
int main() {
//获取CPU芯片数量
int num_cpus = Get_CPU_Number_Linux();
//获取GPU芯片数量
int num_gpus = Get_GPU_Number();
//如果没有GPU,则直接返回
if (num_gpus < 1) {
printf("no CUDA capable devices were detected!\n");
return 1;
}
//打印CPU数量和GPU数量
printf("number of host CPUs:\t%d\n", num_cpus);
printf("number of device GPUs:\t%d\n", num_gpus);
//打印每一个GPU设备的信息
for (int i = 0; i < num_gpus; i++) {
cudaDeviceProp dprop;
cudaGetDeviceProperties(&dprop, i);
printf(" %d: %s\n", i, dprop.name);
}
printf("---------------------------\n");
/
// 初始化数据
int n = num_gpus * 8192;
int nbytes = n * sizeof(int);
int *data = 0; // CPU数据指针
int b = 3; // 数组单个元素增加值
data = (int *)malloc(nbytes); //CPU内数组分配内存
//检查CPU是否成功分配内存
if (0 == data) {
printf("couldn't allocate CPU memory\n");
return 1;
}
//对CPU里的数组数据进行赋值
for (int i = 0; i < n; i++) data[i] = i;
printf("---------------------------\n");
int thread_num = num_cpus;
int* idx = (int*) malloc(thread_num*sizeof(int));
finish_code = (int*) malloc(thread_num*sizeof(int));
for (int i=0;i<thread_num;i++){
pthread_t thread;
idx[i] = i;
finish_code[i] = 1;
address* addr = (address*) malloc(5*sizeof(address));
addr[0] = (address)(&(idx[i]));
addr[1] = (address)(&(n));
addr[2] = (address)(&(nbytes));
addr[3] = (address)(&(b));
addr[4] = (address)(data);
pthread_create(&thread,NULL,exec,addr);
pthread_join(thread, NULL);
}
for (int i=0;i<thread_num;i++){
//while (finish_code[i]) int a = 0;
}
//检查结果是否正确
bool bResult = correctResult(data, n, b);
if (data) free(data); // 释放CPU资源
if (bResult) printf("SUCCESS!\n");
if (!bResult) printf("FAIL!\n");
return 0;
}