symforce.caspar.memory.accessors module

class Common(accessor, idx)[source]

Bases: object

Helper class to make code templating easier.

Parameters:
class DtypeKwargs[source]

Bases: TypedDict

Helper class for dtype-related kwargs.

dtype: DType
kernel_dtype: DType
class Accessor(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: object

Parent class for all memory accessors.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]]
KERNEL_INIT_LINES_TEMPLATE: ClassVar[str] = ''
PY_SIG_TEMPLATE: ClassVar[dict[str, str]]
PY_ARGS_TEMPLATE: ClassVar[list[str]]
SKIP_IF_ALL_ZERO: ClassVar[bool] = False
get_signature(name)[source]
Parameters:

name (str)

Return type:

dict[str, str]

get_layout(storage)[source]
Parameters:

storage (Any)

Return type:

list[list[int]]

pre_kernel_code()[source]
Return type:

str

pre_calc_code()[source]
Return type:

str

post_calc_code()[source]
Return type:

str

shared_size_req()[source]
Return type:

int

stride_string()[source]
Return type:

str

offset_string()[source]
Return type:

str

class ReadSequential(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Sequential, _ReadAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class WriteSequential(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Sequential, _WriteAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class ReadStrided(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Strided, _ReadAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class ReadStridedWithDefault(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Strided, _ReadAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class WriteStrided(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Strided, _WriteAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class AddSequential(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Sequential, _AddAccessor

Accessor for sequentially adding to the output.

Each thread reads, increments and writes to the element at its global thread index.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class ReadIndexed(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Indexed, _ReadAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class WriteIndexed(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Indexed, _WriteAccessor

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class AddIndexed(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Indexed, _AddAccessor

Accessor for adding to indexed elements.

Each thread reads, increments and writes to the index specified by the input array. This accessor does not use atomic operations, so the indices have to be unique.

Not optimized for shared memory access or coalescing.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class ReadShared(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _UsingIndexData, _UsingSharedMem, _ReadAccessor

Accessor for shared memory read access.

You need to generate the shared indices using the lib.shared_indices function. All reads within a block are sorted (for better coalescence), read once and distributed within the block. There is a small overhead compared to Indexed access.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{idx_name}_indices': 'SharedIndex*', '{name}': '{storage_t}*', '{name}_num_alloc': 'unsigned int'}
KERNEL_INIT_LINES_TEMPLATE: ClassVar[str] = '__shared__ SharedIndex {idx_name}_indices_loc[{block_size}];\n{idx_name}_indices_loc[threadIdx.x] = (global_thread_idx<problem_size ?\n                               {idx_name}_indices[global_thread_idx] :\n                               SharedIndex{{0xffffffff, 0xffff, 0xffff}});'
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{idx_name}_indices': 'pybind11::object', '{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})', 'reinterpret_cast<SharedIndex*>(AsUint2Ptr({idx_name}_indices))']
read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class ReadUnique(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _UsingSharedMem, _ReadAccessor

Accessor for shared memory read access. You need to generate the shared indices using the lib.shared_indices function. All reads within a block are sorted (for better coalesence), read once and distributed within the block. There is a small overhead compared to Indexed access.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'const {storage_t}* const'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})']
read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class AddSharedSum(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _UsingIndexData, _UsingSharedMem, _AddAccessor

Accessor for shared sum memory write access.

Each thread adds to the value at a give index. You need to generate the shared indices from the indices using the lib.shared_indices function. All writes within a block are sorted (for better coalesence), written once and distributed within the block.

Equivalent to: (for i, k in enumerate(indices)): out[k] += values[i]

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{idx_name}_indices': 'SharedIndex*', '{name}': '{storage_t}* const', '{name}_num_alloc': 'unsigned int'}
KERNEL_INIT_LINES_TEMPLATE: ClassVar[str] = '__shared__ SharedIndex {idx_name}_indices_loc[{block_size}];\n{idx_name}_indices_loc[threadIdx.x] = (global_thread_idx<problem_size ?\n                               {idx_name}_indices[global_thread_idx] :\n                               SharedIndex{{0xffffffff, 0xffff, 0xffff}});'
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{idx_name}_indices': 'pybind11::object', '{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})', 'reinterpret_cast<SharedIndex*>(AsUint2Ptr({idx_name}_indices))']
write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class WriteBlockSum(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _UsingSharedMem, _WriteAccessor

Accessor for summation over the block.

Each block writes to one element.

To do a full reduction, the user needs to calculate the final sum from the n // 1024 elements.

This class does not use atomic add when writing to the output. You need to generate the shared indices using the lib.shared_indices function.

Equivalent to: (for i, k in enumerate(indices)): out[(k//1024)] += values[i]

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}* const', '{name}_num_alloc': 'unsigned int'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})']
write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class AddBlockSum(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _UsingSharedMem, _AddAccessor

Accessor for summation over the block.

Each block adds to one element.

To do a full reduction, the user needs to calculate the final sum from the n // 1024 elements.

This class does not use atomic add when writing to the output. You need to generate the shared indices using the lib.shared_indices function.

Equivalent to: (for i, k in enumerate(indices)): out[(k//1024)] += values[i]

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}*', '{name}_num_alloc': 'unsigned int'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})']
write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class AddSum(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _UsingSharedMem, _AddAccessor

Accessor for adding global sum.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}* const'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})']
EXTRA_DATA = -992
write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

get_layout(storage)[source]
Parameters:

storage (Any)

Return type:

list[list[int]]

pre_calc_code()[source]
Return type:

str

post_calc_code()[source]
Return type:

str

class WriteSum(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: AddSum

Accessor for writing global sum.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

pre_kernel_code()[source]
Return type:

str

class ReadPair(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Pairwise, _ReadAccessor

Accessor to read the element corresponding to the current thread and the next thread.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}* const', '{name}_num_alloc': 'unsigned int'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})']
read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class ReadPairStridedWithDefault(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Pairwise, _ReadAccessor

Accessor to read the element corresponding to the current thread and the next thread.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}* const', '{name}_num_alloc': 'unsigned int', '{name}_offset': 'int', '{name}_stride': 'int'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object', '{name}_offset': 'int', '{name}_stride': 'int'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})', '{name}_offset', '{name}_stride']
read_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class AddPair(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Pairwise, _AddAccessor

Accessor for adding the pair to the element corresponding to the current thread and the next thread.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}* const', '{name}_num_alloc': 'unsigned int'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})']
EXTRA_DATA = 1
write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class WritePair(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: _Pairwise, _ReadAccessor

Accessor for writing the pair to the element corresponding to the current thread and the next thread. The second element is added to the first element of the next thread.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

KERNEL_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': '{storage_t}* const', '{name}_num_alloc': 'unsigned int'}
PY_SIG_TEMPLATE: ClassVar[dict[str, str]] = {'{name}': 'pybind11::object'}
PY_ARGS_TEMPLATE: ClassVar[list[str]] = ['As{storage_t_capitalized}Ptr({name})', 'GetNumCols({name})']
EXTRA_DATA = 1
pre_kernel_code()[source]
Return type:

str

write_template(idx)[source]
Parameters:

idx (int)

Return type:

str

class TunableShared(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadShared, _TunableAccessor

Used in factors to define a tunable parameter that is shared between factors.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class Tunable(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: TunableShared

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class TunableUnique(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadUnique, _TunableAccessor

Used in factors to define a tunable parameter that is shared between all factors.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class TunablePair(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadPair, _TunableAccessor

Used in factors to define a tunable pair. That is factor[0] depends on arg[0] and arg[1], factor[1] depends on arg[1] and arg[2], etc.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class ConstantSequential(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadSequential, _ConstAccessor

Used in factors to define constants that are unique to each factor.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class Constant(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ConstantSequential

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class ConstantShared(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadShared, _ConstAccessor

Used in factors to define constants that are shared between factors.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class ConstantIndexed(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadIndexed, _ConstAccessor

Used in factors to define constants that are indexed.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)

class ConstantUnique(name, storage, *, dtype, kernel_dtype, reuse_indices_from=None, offset=None, stride=None, default=None, after=None, block_size=1024)[source]

Bases: ReadUnique, _ConstAccessor

Used in factors to define a constant that is shared between all factors.

Parameters:

  • name (str)

  • storage (T.StorableOrType)

  • dtype (DType)

  • kernel_dtype (DType)

  • reuse_indices_from (str | None)

  • offset (int | None)

  • stride (int | None)

  • default (T.Any | None)

  • after (str | None)

  • block_size (int)