Commit 235b9062 authored by Martin Bauer's avatar Martin Bauer
Browse files

lbmpy: cuda & square channel scenario

parent 98176304
......@@ -12,7 +12,7 @@ import hashlib
from pystencils.transformations import symbolNameToVariableName
CONFIG_GCC = {
'compiler': 'g++',
'compiler': 'g++-4.8',
'flags': '-Ofast -DNDEBUG -fPIC -shared -march=native -fopenmp',
}
CONFIG_INTEL = {
......
......@@ -43,7 +43,6 @@ def createKernel(listOfEquations, functionName="kernel", typeForSymbol=None, spl
fieldsRead, fieldsWritten, assignments = typeAllEquations(listOfEquations, typeForSymbol)
allFields = fieldsRead.union(fieldsWritten)
readOnlyFields = set([f.name for f in fieldsRead - fieldsWritten])
body = ast.Block(assignments)
......
......@@ -58,7 +58,7 @@ def buildNumpyArgumentList(kernelFunctionNode, argumentDict):
def makePythonFunction(kernelFunctionNode, argumentDict={}):
mod = SourceModule(str(generateCUDA(kernelFunctionNode)))
mod = SourceModule(str(generateCUDA(kernelFunctionNode)), options=["-w"])
func = mod.get_function(kernelFunctionNode.functionName)
def wrapper(**kwargs):
......
......@@ -25,7 +25,7 @@ def getLinewiseCoordinates(field, ghostLayers):
return arrShape[result.index(cudaIdx)] - 2 * ghostLayers
return {'block': tuple([getShapeOfCudaIdx(idx) for idx in THREAD_IDX]),
'grid': tuple([getShapeOfCudaIdx(idx) for idx in BLOCK_IDX]) }
'grid': tuple([getShapeOfCudaIdx(idx) for idx in BLOCK_IDX])}
return [i + ghostLayers for i in result], getCallParameters
......@@ -37,18 +37,16 @@ def createCUDAKernel(listOfEquations, functionName="kernel", typeForSymbol=None)
typeForSymbol = typingFromSympyInspection(listOfEquations, "float")
fieldsRead, fieldsWritten, assignments = typeAllEquations(listOfEquations, typeForSymbol)
readOnlyFields = set([f.name for f in fieldsRead - fieldsWritten])
allFields = fieldsRead.union(fieldsWritten)
readOnlyFields = set([f.name for f in fieldsRead - fieldsWritten])
code = KernelFunction(Block(assignments), fieldsRead.union(fieldsWritten), functionName)
code = KernelFunction(Block(assignments), allFields, functionName)
code.globalVariables.update(BLOCK_IDX + THREAD_IDX)
fieldAccesses = code.atoms(Field.Access)
requiredGhostLayers = max([fa.requiredGhostLayers for fa in fieldAccesses])
coordMapping, getCallParameters = getLinewiseCoordinates(list(fieldsRead)[0], requiredGhostLayers)
allFields = fieldsRead.union(fieldsWritten)
basePointerInfo = [['spatialInner0']]
basePointerInfos = {f.name: parseBasePointerInfo(basePointerInfo, [2, 1, 0], f) for f in allFields}
......
import sympy as sp
import numpy as np
class SliceMaker(object):
......@@ -78,3 +79,22 @@ def sliceFromDirection(directionName, dim, normalOffset=0, tangentialOffset=0):
assert lowName not in directionName, "Invalid direction name"
result[dimIdx] = normalSliceHigh
return tuple(result)
def removeGhostLayers(arr, indexDimensions=0, ghostLayers=1):
dimensions = len(arr.shape)
spatialDimensions = dimensions - indexDimensions
indexing = [slice(ghostLayers, -ghostLayers, None), ] * spatialDimensions
indexing += [slice(None, None, None)] * indexDimensions
return arr[indexing]
def addGhostLayers(arr, indexDimensions=0, ghostLayers=1):
dimensions = len(arr.shape)
spatialDimensions = dimensions - indexDimensions
newShape = [e + 2 * ghostLayers for e in arr.shape[:spatialDimensions]] + list(arr.shape[spatialDimensions:])
result = np.zeros(newShape)
indexing = [slice(ghostLayers, -ghostLayers, None), ] * spatialDimensions
indexing += [slice(None, None, None)] * indexDimensions
result[indexing] = arr
return result
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment