//信号量
int t = 0;
sem_t empty,full;//信号量 empty表示空的个数加一 full表示满的个数减一
pthread_mutex_t mutex; // 互斥量
void* producer(void* arg){
int* time=(int*) arg;
while(true){
sem_wait(&empty);//空的个数加一
pthread_mutex_lock(&mutex);
//add
t++;
printf("producer add 1 to %dn", t);
pthread_mutex_unlock(&mutex);
sem_post(&full); //满的个数减一
sleep(*time);
}
}
void* customer(void* arg){
while(true){
sem_wait(&full); //满的个数减一
pthread_mutex_lock(&mutex);
//delete
t--;
printf("customer delete 1 to %dn", t);
pthread_mutex_unlock(&mutex);
sem_post(&empty);//空的个数减一
sleep(4);
}
}
int main(){
pthread_t pthread_producer,pthread_producer_2;
pthread_t pthread_customer;
pthread_mutex_init(&mutex,NULL);
sem_init(&full,0,0);
sem_init(&empty,0,10);
int a = 2;
int b = 3;
pthread_create(&pthread_producer,NULL,producer,&a);
pthread_create(&pthread_producer_2,NULL,producer,&b);
pthread_create(&pthread_customer,NULL,customer,NULL);
pthread_join(pthread_customer,NULL);
pthread_join(pthread_producer,NULL);
pthread_join(pthread_producer_2,NULL);
return 0;
}
//条件变量
int t = 0; // 消费者减t，生产者加t
pthread_mutex_t mutex; // 互斥量
pthread_cond_t cond_full; // 非满条件变量
pthread_cond_t cond_empty; // 非空条件变量
int main(){
pthread_t producer;
pthread_t consumer;
// 初始化互斥量
pthread_mutex_init(&mutex,NULL);
// 初始化条件变量
pthread_cond_init(&cond_full,NULL);
pthread_cond_init(&cond_empty,NULL);
// 创建线程
pthread_create(
&producer, // 指明哪个线程
NULL,
producer_function, // 线程执行什么函数
NULL // 传给函数的参数
);
pthread_create(
&consumer,
NULL,
consumer_function,
NULL
);
// 等待线程结束
pthread_join(producer,NULL);
pthread_join(consumer,NULL);
// 销毁
pthread_mutex_destroy(&mutex);
pthread_cond_destroy(&cond_empty);
pthread_cond_destroy(&cond_full);
return 0;
}
void* producer_function(){
while(true){
pthread_mutex_lock(&mutex); // 加锁
// 当生产10个，认为满了
while(t>10){
// 将mutex解锁
// 当收到pthread_cond_signal(&cond_full)时,mutex加锁并向下运行
pthread_cond_wait(&cond_full,&mutex);
}
t++; // 生产
printf("生产"); // 随便写个printf意思意思
pthread_mutex_unlock(&mutex); // 解锁
pthread_cond_signal(&cond_empty);
sleep(5);
}
}
void* consumer_function(){
while(true){
pthread_mutex_lock(&mutex);
while(t==0){
pthread_cond_wait(&cond_empty,&mutex);
}
t--;
printf("消费");
pthread_mutex_unlock(&mutex);
pthread_cond_signal(&cond_full);
sleep(5);
}
}
// 简单并行规约
__global__ void reduceNeighbored(int *d_data, int *d_local_sum, int N)
{
int tid = threadIdx.x;
int index = blockIdx.x * blockDim.x + threadIdx.x;
int *data = d_data + blodckIdx.x * blockDim.x;
if(index >= N) return;
for(int strize = 1; strize < blockDim.x; strize *= 2)
{
if(tid % (2 * strize) == 0)
data[tid] += data[tid + strize];
__synthreads();
}
if(tid == 0) d_local_sum[blockIdx.x] = data[0];
}
// 计算e值
long long num = 100;
int thread_count = 32;
pthread_mutex_t mutex;
double e = 0;
void *thread_function(void *rank)
{
double result = 0;
int rank = (int *)rank;
int cur_n = (num + thread_count-1)/thread_count;
int first = cur_n*rank;
int last = first + cur_n;
for(int i = first; i < last; i++){
result += 1/(double*)f(i); // f(i)为i的阶乘
}
pthread_mutex_lock(&mutex);
e += result;
pthread_mutex_unlock(&mutex);
return null;
}
int main(){
int value[thread_count];
pthread_t thread_id[thread_count];
for(int i = 0; i < thread_count; i++)
{
value[i] = i;
}
for(int i = 0; i < thread_count; i++)
{
pthread_create(&thread_id[i],null, thread_function,(void*) &value[i]);
}
for(int i = 0; i < thread_count; i++){
pthread_join(thread_id[i], null);
}
printf(e);
}
// 矢量求和
#define N (2048*2048)
#define blocksize 256
#define gridsize 512
int main()
{
int *a, *b, *c;
int *d_a, *d_b, *d_c;
int size = N * sizeof(int);
// 为GPU和数据分配空间
cudaMalloc((void**) &d_a, size);
cudaMalloc((void**) &d_b, size);
cudaMalloc((void**) &d_c, size);
a = (int*) malloc(size);
b = (int*) malloc(size);
c = (int*) malloc(size);
for(int i = 0; i < N; i++)
{
a[i] = 1; b[i] = 2;
}
// 内存中的数据放到显存
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
add<<>>(d_a, d_b, d_c);
//显存中数据取回内存
cudamemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
free(a);free(b);free(c);
cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);
return 0;
}
//核函数
__global__ cuda_add(int *a,int *b, int *c){
int index = blockDim.x * blockIdx.x + threadIdx.x;
int thread_sum = blockDim.x * gridDim.x;// 网格中线程总数 因为线程数量小于数据数量 线程需要重复使用
for(int i = index; i < N; i += thread_sum)
{
c[i] = a[i]+b[i];
}
}
// 矩阵转置
const int TILE_DIM = 32;
const int N = 32*32;
const int grid_size_x = (N + TILE_DIM - 1)/TILE_DIM;
const int grid_size_y = grid_size_x;
const dim3 block_size(TILE_DIM, TILEDIM);
const dim3 grid_size(grid_size_x, grid_size_y);
int main()
{
int *a, *b;
int *d_a, *d_b;
int size = N*sizeof(int)
a = (int*) malloc(size);
b = (int*) malloc(size);
cudaMalloc((void**) &d_a, size);
cudaMalloc((void**) &d_b, size);
for(int i = 0; i < TILE_DIM; i++){
for(int j = 0; j < TILE_DIM; j++){
a[i][j] = i+j;
}
}
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cuda_function<<>>(d_a, d_b, TILE_DIM);
cudaMemcpy(b,d_b,size,cudaMemcpyDeviceToHost);
print(*b);
free(a);free(b);
cudaFree(d_a);cudaFree(d_b);
return 0;
}
//普通做法 N 代表一行有多少
__global__ void cuda_function(int *a,int *b, int N)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
int j = blockDim.y * blockIdx.y + threadIdx.y;
if( i < N && j < N)
{
b[j*N+i] = a[i*N+j];
}
}
// 共享内存做法
__global__ void cuda_function(int *a,int *b, int N)
{
__shared__ int s[TILE_DIM][TILE_DIM];
int bx = blockIdx.x * TILE_DIM;
int by = blockIdy.y * TILE_DIM;
int i = bx + threadIdx.x;
int j = by + threadIdx.y;
if(i < N && j < N) // 由直线型转成 转置后并且是二维的
{
s[threadIdx.y][threadIdx.x] = a[j*N+i];
}
__syncThreads();
int i1 = bx + threadIdx.y;
int j1 = by + threadIdx.x;
if(i1 < N && j1 < N)
{
b[i1 * N + j1] = s[thread.x][thread.y];
}
}
//并行规约
int main()
{
const int N = 1 << 29; //512M
int local_length = 1024;
long long total_sum = 0;
dim3 grid(((N + local_length -1) / local_length), 1);
dim3 block(local_length, 1);
int *h_data = null;
int *h_local_num = null;
int *d_data = null;
int *d_local_sum = null;
int size = N * sizeof(int);
h_data = (int*) malloc(size);
h_local_sum = (int*) malloc(int(grid.x)* sizeof(int));
cudaMalloc((void**) &d_data, size);
cudaMalloc((void**) &d_local_sum, int(grid.x) * sizeof(int));
for(int i = 0; i < N; i++)
h_data[i] = int(10 * sin(0.02*3.14*i));
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
reduceNeighbored<<>>(d_data, d_local,N);
cudaMemcpy(h_local_sum, d_data_sum, int(grid.x)*sizeof(int),
cudaMemcpyDeviceToHost);
// 同步
cudaDeviceSynchronize();
for(int i = 0; i < int(grid.x); i++)
{
total_sum += h_local_sum[i];
}
printf(total_sum);
return 0;
}
// 相邻配对
__global__ void reduceNeighbored(int *d_data, int *d_local_sum, int N)
{
int index = blockDim.x * blockIdx.x + threadIdx.x;
int *data = d_data + blockIdx.x * blockDim.x;
if(index > N) return;
for(int stride = 1; stride < blockDim.x; stride *= 2)
{
if(tid%(stride*2) == 0){
data[tid] += data[tid+stide];
}
__syncThreads();
}
if(index == 0) d_local_num[blockIdx.x] = data[0];
}
//交错配对
__global__ void Kenel(int *d_data, int *d_local_sum, int N)
{
int tid = threadIdx.x;
int index = blockIdx.x * blockDim.x + threadIdx.x;
int *data = d_data + blockIdx.x * blockDim.x;
if(index >= N) return;
for(int stride = blockDim.x/2; stride > 0; stride /= 2)
{
if(tid < stride)
{
data[tid] += data[tid + stride]
}
__syncThreads();
}
if(tid == 0){
d_local_sum[blockIdx.x] = data[0];
}
}
// 矩阵相乘
#define TILE 32 // 块的边长
int width = 256; // 矩阵边长
int n = width * width;
int size = n * sizeof(float);
dim3 gird_size(width/TILE, width/TILE);
dim3 block_size(TILE, TILE);
void main()
{
float *a, *b *c;
float *d_a, *d_b, *d_c;
a = (float*) malloc(size);
b = (float*) malloc(size);
c = (float*) malloc(size);
cudaMalloc((void**) &d_a, size);
cudaMalloc((void**) &d_b, size);
cudaMalloc((void**) &d_c, size);
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
MatrixMulKernel<<>>(d_a, d_b, d_c, width);
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
printf(c);
free...
}
// 普通版本
__global__ void MatrixMulKernel(float *d_a, float *d_b, float *d_c, int width)
{
int row = blockDim.y * blockIdx.y + threadIdx.y;
int cow = blockDim.x * blockIdx.x + threadIdx.x;
//结果放在row行 col列
for(int i = 0; i < width; i++)
{
c[row*width + col] += a[row * width + i] * b[k * width + col];
}
}
// 共享内存版本
__global__ void cuda_function(float *a,float *b, float *c.int width)
{
// 共享内存存储于block中 a,b的数据
// 可以使该block中thread可以快速的读入数据
__shared__ float s_a[TILE][TILE]; // TILE是共享内存的大小
__shared__ float s_b[TILE][TILE];
// s_a, s_b的索引
int s_row = threadIdx.y;
int s_col = threadIdx.x;
// a的row行 b的col乘积的和 放在 c的row行 col列 c[row*width+col]
int row = blockIdx.y * blockDim.x + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
// 遍历grid中的某一行或某一列 block, k代表第k个block
for(int k = 0; k < width/TILE; i++){
//将计算的结果先放在共享内存s_a, s_b
int k_col = k*TILE + threadIdx.x;
int k_row = K*TILE + threadIdx.y;
s_a[s_row][s_col] = a[row * width + k_col];
s_b[s_row][s_col] = b[k_row * width + col];
//等待线程同步
__syncThreads();
float sum = 0;
for(int i = 0; i < TILE; i++)
{
sum += s_a[s_row][i] * s_b[i][s_col];
}
// 等待block内所有线程计算完毕
__syncThreads();
c[row * width + col] = sum;
}
}
