<s> from numba import cuda <EOL> from numba . cuda . cudadrv . driver import driver <EOL> from numba import numpy_support as nps <EOL> import math <EOL> def transpose ( a , b = None ) : <EOL> """<STR_LIT>""" <EOL> if not b : <EOL> cols , rows = a . shape <EOL> strides = a . dtype . itemsize * cols , a . dtype . itemsize <EOL> b = cuda . cudadrv . devicearray . DeviceNDArray ( ( rows , cols ) , strides , dtype = a . dtype ) <EOL> dt = nps . from_dtype ( a . dtype ) <EOL> tpb = driver . get_device ( ) . MAX_THREADS_PER_BLOCK <EOL> tile_width = int ( math . pow ( <NUM_LIT:2> , math . log ( tpb , <NUM_LIT:2> ) / <NUM_LIT:2> ) ) <EOL> tile_height = int ( tpb / tile_width ) <EOL> tile_shape = ( tile_height , tile_width + <NUM_LIT:1> ) <EOL> @ cuda . jit <EOL> def kernel ( input , output ) : <EOL> tile = cuda . shared . array ( shape = tile_shape , dtype = dt ) <EOL> tx = cuda . threadIdx . x <EOL> ty = cuda . threadIdx . y <EOL> bx = cuda . blockIdx . x * cuda . blockDim . x <EOL> by = cuda . blockIdx . y * cuda . blockDim . y <EOL> x = by + tx <EOL> y = bx + ty <EOL> if by + ty < input . shape [ <NUM_LIT:0> ] and bx + tx < input . shape [ <NUM_LIT:1> ] : <EOL> tile [ ty , tx ] = input [ by + ty , bx + tx ] <EOL> cuda . syncthreads ( ) <EOL> if y < output . shape [ <NUM_LIT:0> ] and x < output . shape [ <NUM_LIT:1> ] : <EOL> output [ y , x ] = tile [ tx , ty ] <EOL> blocks = int ( b . shape [ <NUM_LIT:0> ] / tile_height + <NUM_LIT:1> ) , int ( b . shape [ <NUM_LIT:1> ] / tile_width + <NUM_LIT:1> ) <EOL> threads = tile_height , tile_width <EOL> kernel [ blocks , threads ] ( a , b ) <EOL> return b </s>
