Advertisement
Not a member of Pastebin yet?
Sign Up,
it unlocks many cool features!
- #include "cuda_runtime.h"
- #include "device_launch_parameters.h"
- #include <stdio.h>
- const int Nbl = 2;
- const int Nth = 2;
- const int N = 16;
- __global__ void scalar(int *a, int *b, int *res){
- __shared__ int s[Nth];
- int temp = 0;
- int tid=threadIdx.x+blockIdx.x*blockDim.x;
- int ind = threadIdx.x;
- while (tid<N)
- {
- temp+=a[tid]*b[tid];
- tid+=blockDim.x*gridDim.x;
- }
- s[ind]=temp;
- int k = blockDim.x/2;
- __syncthreads();
- while(k>0){
- if(ind<k){
- s[ind] +=s[ind+k];
- }
- __syncthreads();//Функция __syncthreads() обеспечивает синхронизацию потоков в блоке, которая будет ждать до тех пор, пока все запущенные потоки отработают до этой точки. Функция необходима для данных обрабатываемых одним потоком, затем будут использоваться другими потоками.
- k=k/2;
- }
- res[blockIdx.x] = s[0];
- }
- int main()
- {
- int *dev_a, *dev_b, *dev_res;
- int a[N], b[N], res[Nbl];
- for (long i = 0; i < N; i++)
- {
- a[i] = i;
- b[i] = i+1;
- }
- a[5]=-4;
- cudaMalloc((void**)&dev_a, sizeof(int)*N);
- cudaMalloc((void**)&dev_b, sizeof(int)*N);
- cudaMalloc((void**)&dev_res, sizeof(int)*Nbl);
- cudaMemcpy(dev_a, a, sizeof(int)*N, cudaMemcpyHostToDevice);
- cudaMemcpy(dev_b, b, sizeof(int)*N, cudaMemcpyHostToDevice);
- scalar<<<Nbl, Nth >>>(dev_a, dev_b, dev_res);
- cudaMemcpy(res, dev_res, sizeof(int)*Nbl, cudaMemcpyDeviceToHost);
- int s = 0;
- for (int i =0;i<Nbl;i++)
- s+= res[i];
- printf("%d", s);
- printf("\n");
- s = 0;
- for (int i=0;i<N;i++)
- s+=a[i]*b[i];
- printf("%d", s);
- getchar();
- return 0;
- }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement