|
Brick Library 0.1
Performance-portable stencil datalayout & codegen
|
Public Member Functions | |
| def | __init__ (self, VECLEN=32, LID=None, ocl=False) |
| def | checkConfig (self) |
| def | setCodeGen (self, codegen) |
| def | genVectorLoop (self, CodeBlock group) |
| def | gen_lhs (self, Buffer buf, List[int] offset, rel=None, dimrels=None) |
| def | gen_rhs (self, Expr comp, List[int] shift, List[int] offset, rel=None, dimrels=None) |
| def | declare_reg (self, name, CodeBlock block) |
| def | declare_vec (self, name, CodeBlock block) |
| def | declare_gridref (self, Grid grid, CodeBlock block) |
| def | store_vecbuf (self, vecbuf_name, reg_name, CodeBlock block) |
| def | declare_buf (self, Buffer buf, CodeBlock block) |
| def | genStoreLoop (self, CodeBlock group) |
| def | storeTile (self, Buffer buf, CodeBlock group) |
| def | genStoreTileLoop (self, CodeBlock group, dims) |
| def | store (self, Buffer buf, CodeBlock group) |
| def | read_aligned (self, Grid grid, offset, str name, CodeBlock block, rel=None) |
| def | merge (self, rego, regl, regr, dim, shift, CodeBlock block) |
Public Member Functions inherited from st.codegen.backend.base.Backend | |
| def | __init__ (self) |
| def | setCodeGen (self, codegen) |
| def | setLayout (self, layout) |
| def | prequel (self, toplevel) |
| def | checkConfig (self) |
| def | declare_buf (self, Buffer buf, CodeBlock block) |
| def | declare_gridref (self, Grid grid, CodeBlock block) |
| def | stride (self, dim) |
| def | declare_vecbuf (self, Grid grid, vec_shift, CodeBlock block) |
| def | genVectorLoop (self, CodeBlock group) |
| def | genStoreLoop (self, CodeBlock group) |
| def | store (self, Buffer buf, CodeBlock group) |
| def | storeTile (self, Buffer buf, CodeBlock group) |
| def | genStoreLoc (self, Grid grid, shift, offset, rel, dimrels) |
| def | genStoreTileLoop (self, CodeBlock group, dims) |
| def | gen_lhs (self, Buffer buf, List[int] offset, rel=None, dimrels=None) |
| def | gen_rhs (self, Expr comp, List[int] shift, List[int] offset, rel=None, dimrels=None) |
| def | declare_reg (self, name, CodeBlock block) |
| def | declare_vec (self, name, CodeBlock block) |
| def | store_vecbuf (self, vecbuf_name, reg_name, CodeBlock block) |
| def | merge (self, rego, regl, regr, dim, shift, CodeBlock block) |
| def | read_aligned (self, Grid grid, offset, str name, CodeBlock block, rel=None) |
Public Attributes | |
| VECLEN | |
| ocl | |
| LID | |
| STRIDE | |
| codegen | |
| printer | |
Public Attributes inherited from st.codegen.backend.base.Backend | |
| printer | |
| codegen | |
| prec | |
| layout | |
| STRIDE | |
| VECLEN | |
| ALIGNED | |
Static Public Attributes | |
| str | LID = "threadIdx.x" |
Static Public Attributes inherited from st.codegen.backend.base.Backend | |
| str | preffix = "_cg" |
Additional Inherited Members | |
Static Public Member Functions inherited from st.codegen.backend.base.Backend | |
| def | gridref_name (grid) |
| def | vectmp_name (idx) |
| def | vecbuf_name (grid, vec_shift) |
| def | vecreg_name (grid, vec_shift) |
| def | index_name (idx) |
| def | rel_name (idx=None) |
| def st.codegen.backend.cuda.BackendCUDA.__init__ | ( | self, | |
VECLEN = 32, |
|||
LID = None, |
|||
ocl = False |
|||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
Reimplemented in st.codegen.backend.cuda.BackendCuFlex.
| def st.codegen.backend.cuda.BackendCUDA.checkConfig | ( | self | ) |
Reimplemented from st.codegen.backend.base.Backend.
Reimplemented in st.codegen.backend.cuda.BackendCuFlex.
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.declare_reg | ( | self, | |
| name, | |||
| CodeBlock | block | ||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.declare_vec | ( | self, | |
| name, | |||
| CodeBlock | block | ||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
Reimplemented in st.codegen.backend.cuda.BackendCuFlex.
| def st.codegen.backend.cuda.BackendCUDA.gen_lhs | ( | self, | |
| Buffer | buf, | ||
| List[int] | offset, | ||
rel = None, |
|||
dimrels = None |
|||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.gen_rhs | ( | self, | |
| Expr | comp, | ||
| List[int] | shift, | ||
| List[int] | offset, | ||
rel = None, |
|||
dimrels = None |
|||
| ) |
:param comp: The expression to print :param shift: The shift of scatter :param offset: Scattered from :param rel: Added offset when using loops :return:
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.genStoreLoop | ( | self, | |
| CodeBlock | group | ||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.genStoreTileLoop | ( | self, | |
| CodeBlock | group, | ||
| dims | |||
| ) |
| def st.codegen.backend.cuda.BackendCUDA.genVectorLoop | ( | self, | |
| CodeBlock | group | ||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.merge | ( | self, | |
| rego, | |||
| regl, | |||
| regr, | |||
| dim, | |||
| shift, | |||
| CodeBlock | block | ||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
Reimplemented in st.codegen.backend.cuda.BackendCuFlex.
| def st.codegen.backend.cuda.BackendCUDA.read_aligned | ( | self, | |
| Grid | grid, | ||
| offset, | |||
| str | name, | ||
| CodeBlock | block, | ||
rel = None |
|||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
Reimplemented in st.codegen.backend.cuda.BackendCuFlex.
| def st.codegen.backend.cuda.BackendCUDA.setCodeGen | ( | self, | |
| codegen | |||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
| def st.codegen.backend.cuda.BackendCUDA.store_vecbuf | ( | self, | |
| vecbuf_name, | |||
| reg_name, | |||
| CodeBlock | block | ||
| ) |
Reimplemented from st.codegen.backend.base.Backend.
Reimplemented in st.codegen.backend.cuda.BackendCuFlex.
| st.codegen.backend.cuda.BackendCUDA.codegen |
|
static |
| st.codegen.backend.cuda.BackendCUDA.LID |
| st.codegen.backend.cuda.BackendCUDA.ocl |
| st.codegen.backend.cuda.BackendCUDA.printer |
| st.codegen.backend.cuda.BackendCUDA.STRIDE |
| st.codegen.backend.cuda.BackendCUDA.VECLEN |