Blame view

stim/cuda/arraymath/array_cos.cuh 1.44 KB
f186dbda   Tianshu Cheng   header file for b...
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
  #ifndef STIM_CUDA_ARRAY_COS_H
  #define STIM_CUDA_ARRAY_COS_H
  
  #include <iostream>
  #include <cuda.h>
  #include <cmath>
  #include <stim/cuda/cudatools.h>
  
  namespace stim{
  	namespace cuda{
  
  		template<typename T>
  		__global__ void cuda_cos(T* ptr1, T* out, unsigned int N){
  
  			//calculate the 1D index for this thread
  			int idx = blockIdx.x * blockDim.x + threadIdx.x;
  
  			if(idx < N){
  				out[idx] = cos(ptr1[idx]);
  			}
  
  		}
  
  		template<typename T>
  		void gpu_cos(T* ptr1, T* out, unsigned int N){
  
  			//get the maximum number of threads per block for the CUDA device
  			int threads = stim::maxThreadsPerBlock();
  
  			//calculate the number of blocks
  			int blocks = N / threads + 1;
  
  			//call the kernel to do the multiplication
  			cuda_cos <<< blocks, threads >>>(ptr1, out, N);
  
  		}
  
  		template<typename T>
  		void cpu_cos(T* ptr1, T* cpu_out, unsigned int N){
  
  			//allocate memory on the GPU for the array
  			T* gpu_ptr1; 
  			T* gpu_out;
  			HANDLE_ERROR( cudaMalloc( &gpu_ptr1, N * sizeof(T) ) );
  			HANDLE_ERROR( cudaMalloc( &gpu_out, N * sizeof(T) ) );
  
  			//copy the array to the GPU
  			HANDLE_ERROR( cudaMemcpy( gpu_ptr1, ptr1, N * sizeof(T), cudaMemcpyHostToDevice) );
  
  			//call the GPU version of this function
  			gpu_cos<T>(gpu_ptr1 ,gpu_out, N);
  
  			//copy the array back to the CPU
  			HANDLE_ERROR( cudaMemcpy( cpu_out, gpu_out, N * sizeof(T), cudaMemcpyDeviceToHost) );
  
  			//free allocated memory
  			cudaFree(gpu_ptr1);
  			cudaFree(gpu_out);
  
  		}
  		
  	}
  }
  
  
  
  #endif