numbapig.py 3.11 KB
 Neil Gershenfeld committed Feb 09, 2020 1 2 3 ``````# # numbapig.py # Neil Gershenfeld 2/9/20 `````` Neil Gershenfeld committed Feb 09, 2020 4 ``````# calculation of pi by a Numba GPU sum `````` Neil Gershenfeld committed Feb 09, 2020 5 6 7 8 9 10 11 12 13 ``````# pi = 3.14159265358979323846 # from numba import cuda import numpy as np import time # # problem size # block_size = 2**10 `````` Neil Gershenfeld committed Feb 09, 2020 14 ``````grid_size = 2**21 `````` Neil Gershenfeld committed Feb 09, 2020 15 16 ``````NPTS = grid_size*block_size # `````` Neil Gershenfeld committed Feb 09, 2020 17 ``````# kernels and functions `````` Neil Gershenfeld committed Feb 09, 2020 18 19 20 21 ``````# @cuda.jit def init(arr): i = 1+cuda.grid(1) `````` Neil Gershenfeld committed Feb 09, 2020 22 `````` arr[i-1] = 0.5/((i-0.75)*(i-0.25)) `````` Neil Gershenfeld committed Feb 09, 2020 23 `````` #arr[i-1] = i # for testing reduction `````` Neil Gershenfeld committed Feb 09, 2020 24 ``````# `````` Neil Gershenfeld committed Feb 09, 2020 25 ``````@cuda.reduce `````` Neil Gershenfeld committed Feb 09, 2020 26 ``````def Numba_reduce(a,b): `````` Neil Gershenfeld committed Feb 09, 2020 27 28 `````` return a+b # `````` Neil Gershenfeld committed Feb 09, 2020 29 30 31 32 33 34 35 36 37 38 39 40 41 42 ``````@cuda.jit def CUDA_sum(arr,len): i = cuda.grid(1) if (i < len): arr[i] += arr[i+len] # def CUDA_reduce(arr,NPTS): len = NPTS >> 1 while (1): CUDA_sum[grid_size,block_size](arr,len) len = len >> 1 if (len == 0): return # `````` Neil Gershenfeld committed Feb 09, 2020 43 44 45 46 47 48 ``````@cuda.jit def CUDA_result(arr,result): i = cuda.grid(1) if (i == 0): result[0] = arr[0] # `````` Neil Gershenfeld committed Feb 09, 2020 49 ``````# device array `````` Neil Gershenfeld committed Feb 09, 2020 50 51 ``````# arr = cuda.device_array(NPTS,np.float32) `````` Neil Gershenfeld committed Feb 09, 2020 52 53 54 ``````result = cuda.device_array(1,np.float32) #arr = cuda.device_array(NPTS,np.int64) # for testing reduction #result = cuda.device_array(1,np.int64) # for testing reduction `````` Neil Gershenfeld committed Feb 09, 2020 55 ``````# `````` Neil Gershenfeld committed Feb 09, 2020 56 ``````# compile kernels by calling them `````` Neil Gershenfeld committed Feb 09, 2020 57 ``````# `````` Neil Gershenfeld committed Feb 09, 2020 58 ``````init[grid_size,block_size](arr) `````` Neil Gershenfeld committed Feb 09, 2020 59 60 ``````pi = Numba_reduce(arr) CUDA_reduce(arr,NPTS) `````` Neil Gershenfeld committed Feb 09, 2020 61 ``````CUDA_result(arr,result) `````` Neil Gershenfeld committed Feb 09, 2020 62 ``````# `````` Neil Gershenfeld committed Feb 09, 2020 63 ``````# CUDA kernel array calculation `````` Neil Gershenfeld committed Feb 09, 2020 64 65 66 67 68 ``````# start_time = time.time() init[grid_size,block_size](arr) end_time = time.time() mflops = NPTS*4.0/(1.0e6*(end_time-start_time)) `````` Neil Gershenfeld committed Feb 09, 2020 69 ``````print("CUDA kernel array calculation:") `````` Neil Gershenfeld committed Feb 09, 2020 70 71 ``````print(" time = %f, estimated MFlops = %f"%(end_time-start_time,mflops)) # `````` Neil Gershenfeld committed Feb 09, 2020 72 ``````# Numba reduce `````` Neil Gershenfeld committed Feb 09, 2020 73 ``````# `````` Neil Gershenfeld committed Feb 09, 2020 74 ``````init[grid_size,block_size](arr) `````` Neil Gershenfeld committed Feb 09, 2020 75 ``````start_time = time.time() `````` Neil Gershenfeld committed Feb 09, 2020 76 ``````pi = Numba_reduce(arr) `````` Neil Gershenfeld committed Feb 09, 2020 77 78 ``````end_time = time.time() mflops = NPTS*1.0/(1.0e6*(end_time-start_time)) `````` Neil Gershenfeld committed Feb 09, 2020 79 ``````print("Numba reduce:") `````` Neil Gershenfeld committed Feb 09, 2020 80 81 ``````print(" time = %f, estimated MFlops = %f"%(end_time-start_time,mflops)) # `````` Neil Gershenfeld committed Feb 09, 2020 82 ``````# both with Numba reduce `````` Neil Gershenfeld committed Feb 09, 2020 83 84 85 ``````# start_time = time.time() init[grid_size,block_size](arr) `````` Neil Gershenfeld committed Feb 09, 2020 86 ``````pi = Numba_reduce(arr) `````` Neil Gershenfeld committed Feb 09, 2020 87 88 ``````end_time = time.time() mflops = NPTS*5.0/(1.0e6*(end_time-start_time)) `````` Neil Gershenfeld committed Feb 09, 2020 89 ``````print("both with Numba reduce:") `````` Neil Gershenfeld committed Feb 09, 2020 90 91 ``````print(" NPTS = %d, pi = %f"%(NPTS,pi)) print(" time = %f, estimated MFlops = %f"%(end_time-start_time,mflops)) `````` Neil Gershenfeld committed Feb 09, 2020 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 ``````# # CUDA kernel reduction # init[grid_size,block_size](arr) start_time = time.time() CUDA_reduce(arr,NPTS) end_time = time.time() mflops = NPTS*1.0/(1.0e6*(end_time-start_time)) print("CUDA kernel reduction:") print(" time = %f, estimated MFlops = %f"%(end_time-start_time,mflops)) # # both with CUDA kernel reduction # start_time = time.time() init[grid_size,block_size](arr) CUDA_reduce(arr,NPTS) `````` Neil Gershenfeld committed Feb 09, 2020 108 ``````CUDA_result(arr,result) `````` Neil Gershenfeld committed Feb 09, 2020 109 ``````end_time = time.time() `````` Neil Gershenfeld committed Feb 09, 2020 110 ``````pi = result.copy_to_host() `````` Neil Gershenfeld committed Feb 09, 2020 111 112 ``````mflops = NPTS*5.0/(1.0e6*(end_time-start_time)) print("both with CUDA kernel reduction:") `````` Neil Gershenfeld committed Feb 09, 2020 113 ``````print(" NPTS = %d, pi = %f"%(NPTS,pi[0])) `````` Neil Gershenfeld committed Feb 09, 2020 114 ``````print(" time = %f, estimated MFlops = %f"%(end_time-start_time,mflops)) `````` Neil Gershenfeld committed Feb 10, 2020 115 116 117 118 119 120 121 122 123 124 125 126 127 ``````# # both with CUDA kernel reduction and transfer # start_time = time.time() init[grid_size,block_size](arr) CUDA_reduce(arr,NPTS) CUDA_result(arr,result) pi = result.copy_to_host() end_time = time.time() mflops = NPTS*5.0/(1.0e6*(end_time-start_time)) print("both with CUDA kernel reduction and transfer:") print(" NPTS = %d, pi = %f"%(NPTS,pi[0])) print(" time = %f, estimated MFlops = %f"%(end_time-start_time,mflops)) `````` Neil Gershenfeld committed Feb 09, 2020 128