<s> from __future__ import print_function , absolute_import <EOL> from numba import void , float32 , float64 <EOL> import numpy as np <EOL> import numpy . core . umath_tests as ut <EOL> from numba import guvectorize <EOL> from numba import cuda <EOL> from timeit import default_timer as time <EOL> from numba import unittest_support as unittest <EOL> from numba . cuda . testing import skip_on_cudasim <EOL> non_stream_speedups = [ ] <EOL> stream_speedups = [ ] <EOL> @ skip_on_cudasim ( '<STR_LIT>' ) <EOL> class TestCUDAGufunc ( unittest . TestCase ) : <EOL> def test_gufunc_small ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> gufunc . max_blocksize = <NUM_LIT> <EOL> matrix_ct = <NUM_LIT:2> <EOL> A = np . arange ( matrix_ct * <NUM_LIT:2> * <NUM_LIT:4> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:2> , <EOL> <NUM_LIT:4> ) <EOL> B = np . arange ( matrix_ct * <NUM_LIT:4> * <NUM_LIT:5> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:4> , <EOL> <NUM_LIT:5> ) <EOL> ts = time ( ) <EOL> C = gufunc ( A , B ) <EOL> tcuda = time ( ) - ts <EOL> ts = time ( ) <EOL> Gold = ut . matrix_multiply ( A , B ) <EOL> tcpu = time ( ) - ts <EOL> non_stream_speedups . append ( tcpu / tcuda ) <EOL> print ( C , Gold ) <EOL> self . assertTrue ( np . allclose ( C , Gold ) ) <EOL> def test_gufunc_auto_transfer ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> gufunc . max_blocksize = <NUM_LIT> <EOL> matrix_ct = <NUM_LIT:2> <EOL> A = np . arange ( matrix_ct * <NUM_LIT:2> * <NUM_LIT:4> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:2> , <EOL> <NUM_LIT:4> ) <EOL> B = np . arange ( matrix_ct * <NUM_LIT:4> * <NUM_LIT:5> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:4> , <EOL> <NUM_LIT:5> ) <EOL> dB = cuda . to_device ( B ) <EOL> ts = time ( ) <EOL> C = gufunc ( A , dB ) . copy_to_host ( ) <EOL> tcuda = time ( ) - ts <EOL> ts = time ( ) <EOL> Gold = ut . matrix_multiply ( A , B ) <EOL> tcpu = time ( ) - ts <EOL> non_stream_speedups . append ( tcpu / tcuda ) <EOL> print ( C , Gold ) <EOL> self . assertTrue ( np . allclose ( C , Gold ) ) <EOL> def test_gufunc ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> gufunc . max_blocksize = <NUM_LIT> <EOL> matrix_ct = <NUM_LIT> <EOL> A = np . arange ( matrix_ct * <NUM_LIT:2> * <NUM_LIT:4> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:2> , <EOL> <NUM_LIT:4> ) <EOL> B = np . arange ( matrix_ct * <NUM_LIT:4> * <NUM_LIT:5> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:4> , <EOL> <NUM_LIT:5> ) <EOL> ts = time ( ) <EOL> C = gufunc ( A , B ) <EOL> tcuda = time ( ) - ts <EOL> ts = time ( ) <EOL> Gold = ut . matrix_multiply ( A , B ) <EOL> tcpu = time ( ) - ts <EOL> non_stream_speedups . append ( tcpu / tcuda ) <EOL> self . assertTrue ( np . allclose ( C , Gold ) ) <EOL> def test_gufunc_hidim ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> gufunc . max_blocksize = <NUM_LIT> <EOL> matrix_ct = <NUM_LIT:100> <EOL> A = np . arange ( matrix_ct * <NUM_LIT:2> * <NUM_LIT:4> , dtype = np . float32 ) . reshape ( <NUM_LIT:4> , <NUM_LIT> , <NUM_LIT:2> , <NUM_LIT:4> ) <EOL> B = np . arange ( matrix_ct * <NUM_LIT:4> * <NUM_LIT:5> , dtype = np . float32 ) . reshape ( <NUM_LIT:4> , <NUM_LIT> , <NUM_LIT:4> , <NUM_LIT:5> ) <EOL> ts = time ( ) <EOL> C = gufunc ( A , B ) <EOL> tcuda = time ( ) - ts <EOL> ts = time ( ) <EOL> Gold = ut . matrix_multiply ( A , B ) <EOL> tcpu = time ( ) - ts <EOL> non_stream_speedups . append ( tcpu / tcuda ) <EOL> self . assertTrue ( np . allclose ( C , Gold ) ) <EOL> def test_gufunc_new_axis ( self ) : <EOL> @ guvectorize ( [ void ( float64 [ : , : ] , float64 [ : , : ] , float64 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> X = np . random . randn ( <NUM_LIT:10> , <NUM_LIT:3> , <NUM_LIT:3> ) <EOL> Y = np . random . randn ( <NUM_LIT:3> , <NUM_LIT:3> ) <EOL> gold = ut . matrix_multiply ( X , Y ) <EOL> res1 = gufunc ( X , Y ) <EOL> np . testing . assert_allclose ( gold , res1 ) <EOL> res2 = gufunc ( X , np . tile ( Y , ( <NUM_LIT:10> , <NUM_LIT:1> , <NUM_LIT:1> ) ) ) <EOL> np . testing . assert_allclose ( gold , res2 ) <EOL> def test_gufunc_adjust_blocksize ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> gufunc . max_blocksize = <NUM_LIT> <EOL> matrix_ct = <NUM_LIT> <EOL> A = np . arange ( matrix_ct * <NUM_LIT:2> * <NUM_LIT:4> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:2> , <EOL> <NUM_LIT:4> ) <EOL> B = np . arange ( matrix_ct * <NUM_LIT:4> * <NUM_LIT:5> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:4> , <EOL> <NUM_LIT:5> ) <EOL> gufunc . max_blocksize = <NUM_LIT:32> <EOL> C = gufunc ( A , B ) <EOL> Gold = ut . matrix_multiply ( A , B ) <EOL> self . assertTrue ( np . allclose ( C , Gold ) ) <EOL> def test_gufunc_stream ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def matmulcore ( A , B , C ) : <EOL> m , n = A . shape <EOL> n , p = B . shape <EOL> for i in range ( m ) : <EOL> for j in range ( p ) : <EOL> C [ i , j ] = <NUM_LIT:0> <EOL> for k in range ( n ) : <EOL> C [ i , j ] += A [ i , k ] * B [ k , j ] <EOL> gufunc = matmulcore <EOL> gufunc . max_blocksize = <NUM_LIT> <EOL> matrix_ct = <NUM_LIT> <EOL> A = np . arange ( matrix_ct * <NUM_LIT:2> * <NUM_LIT:4> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:2> , <EOL> <NUM_LIT:4> ) <EOL> B = np . arange ( matrix_ct * <NUM_LIT:4> * <NUM_LIT:5> , dtype = np . float32 ) . reshape ( matrix_ct , <NUM_LIT:4> , <EOL> <NUM_LIT:5> ) <EOL> ts = time ( ) <EOL> stream = cuda . stream ( ) <EOL> dA = cuda . to_device ( A , stream ) <EOL> dB = cuda . to_device ( B , stream ) <EOL> dC = cuda . device_array ( shape = ( <NUM_LIT> , <NUM_LIT:2> , <NUM_LIT:5> ) , dtype = A . dtype , stream = stream ) <EOL> dC = gufunc ( dA , dB , out = dC , stream = stream ) <EOL> C = dC . copy_to_host ( stream = stream ) <EOL> stream . synchronize ( ) <EOL> tcuda = time ( ) - ts <EOL> ts = time ( ) <EOL> Gold = ut . matrix_multiply ( A , B ) <EOL> tcpu = time ( ) - ts <EOL> stream_speedups . append ( tcpu / tcuda ) <EOL> self . assertTrue ( np . allclose ( C , Gold ) ) <EOL> def test_copy ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : ] , float32 [ : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def copy ( A , B ) : <EOL> for i in range ( B . size ) : <EOL> B [ i ] = A [ i ] <EOL> A = np . arange ( <NUM_LIT:10> , dtype = np . float32 ) + <NUM_LIT:1> <EOL> B = np . zeros_like ( A ) <EOL> copy ( A , out = B ) <EOL> self . assertTrue ( np . allclose ( A , B ) ) <EOL> def test_copy_odd ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : ] , float32 [ : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def copy ( A , B ) : <EOL> for i in range ( B . size ) : <EOL> B [ i ] = A [ i ] <EOL> A = np . arange ( <NUM_LIT:11> , dtype = np . float32 ) + <NUM_LIT:1> <EOL> B = np . zeros_like ( A ) <EOL> copy ( A , out = B ) <EOL> self . assertTrue ( np . allclose ( A , B ) ) <EOL> def test_copy2d ( self ) : <EOL> @ guvectorize ( [ void ( float32 [ : , : ] , float32 [ : , : ] ) ] , <EOL> '<STR_LIT>' , <EOL> target = '<STR_LIT>' ) <EOL> def copy2d ( A , B ) : <EOL> for x in range ( B . shape [ <NUM_LIT:0> ] ) : <EOL> for y in range ( B . shape [ <NUM_LIT:1> ] ) : <EOL> B [ x , y ] = A [ x , y ] <EOL> A = np . arange ( <NUM_LIT:30> , dtype = np . float32 ) . reshape ( <NUM_LIT:5> , <NUM_LIT:6> ) + <NUM_LIT:1> <EOL> B = np . zeros_like ( A ) <EOL> copy2d ( A , out = B ) <EOL> self . assertTrue ( np . allclose ( A , B ) ) <EOL> def test_nopython_flag ( self ) : <EOL> def foo ( A , B ) : <EOL> pass <EOL> guvectorize ( [ void ( float32 [ : ] , float32 [ : ] ) ] , '<STR_LIT>' , target = '<STR_LIT>' , <EOL> nopython = True ) ( foo ) <EOL> with self . assertRaises ( TypeError ) as raises : <EOL> guvectorize ( [ void ( float32 [ : ] , float32 [ : ] ) ] , '<STR_LIT>' , <EOL> target = '<STR_LIT>' , nopython = False ) ( foo ) <EOL> self . assertEqual ( "<STR_LIT>" , str ( raises . exception ) ) <EOL> def test_invalid_flags ( self ) : <EOL> def foo ( A , B ) : <EOL> pass <EOL> with self . assertRaises ( TypeError ) as raises : <EOL> guvectorize ( [ void ( float32 [ : ] , float32 [ : ] ) ] , '<STR_LIT>' , <EOL> target = '<STR_LIT>' , what1 = True , ever2 = False ) ( foo ) <EOL> head = "<STR_LIT>" <EOL> msg = str ( raises . exception ) <EOL> self . assertEqual ( msg [ : len ( head ) ] , head ) <EOL> items = msg [ len ( head ) : ] . strip ( ) . split ( '<STR_LIT:U+002C>' ) <EOL> items = [ i . strip ( "<STR_LIT>" ) for i in items ] <EOL> self . assertEqual ( set ( [ '<STR_LIT>' , '<STR_LIT>' ] ) , set ( items ) ) <EOL> if __name__ == '<STR_LIT:__main__>' : <EOL> unittest . main ( ) </s>
