tilelang.jit.adapter.wrapper
============================
.. py:module:: tilelang.jit.adapter.wrapper
Attributes
----------
.. autoapisummary::
tilelang.jit.adapter.wrapper.PREDEF_ATTRIBUTE_SET_DYNAMIC_MEMORY
tilelang.jit.adapter.wrapper.PREDEF_ATTRIBUTE_SET_DYNAMIC_MEMORY_HIP
tilelang.jit.adapter.wrapper.PREDEF_INIT_FUNC
tilelang.jit.adapter.wrapper.PREDEF_HOST_FUNC
tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_CREATE_HANDLE
tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_INIT_FUNC
tilelang.jit.adapter.wrapper.L2_PERSISTENT_MAP_RESET_HANDLE
tilelang.jit.adapter.wrapper.TMA_DESC_INIT_FUNC
tilelang.jit.adapter.wrapper.TMA_IM2COL_DESC_INIT_FUNC
tilelang.jit.adapter.wrapper.logger
Classes
-------
.. autoapisummary::
tilelang.jit.adapter.wrapper.BaseWrapper
tilelang.jit.adapter.wrapper.TLCUDASourceWrapper
tilelang.jit.adapter.wrapper.TLHIPSourceWrapper
tilelang.jit.adapter.wrapper.TLCPUSourceWrapper
tilelang.jit.adapter.wrapper.TLMetalSourceWrapper
tilelang.jit.adapter.wrapper.TLWrapper
tilelang.jit.adapter.wrapper.TLPyWrapper
Module Contents
---------------
.. py:data:: PREDEF_ATTRIBUTE_SET_DYNAMIC_MEMORY
:value: Multiline-String
.. raw:: html
Show Value
.. code-block:: python
"""
cudaError_t result_{0} = cudaFuncSetAttribute({0}, cudaFuncAttributeMaxDynamicSharedMemorySize, {1});
if (result_{0} != cudaSuccess) {{
snprintf(error_buf, ERROR_BUF_SIZE, "Failed to set the allowed dynamic shared memory size to %d with error: %s", {1}, cudaGetErrorString(result_{0}));
return -1;
}}
"""
.. raw:: html
Show Value
.. code-block:: python
"""
if ({1} > 65536) {{
snprintf(error_buf, ERROR_BUF_SIZE, "Failed to set the allowed dynamic shared memory size for {0} to %d", {1});
return -1;
}}
return 0;
"""
.. raw:: html
Show Value
.. code-block:: python
"""
#define ERROR_BUF_SIZE 1024
static char error_buf[ERROR_BUF_SIZE];
extern "C" const char* get_last_error() {{
return error_buf;
}}
extern "C" int init() {{
error_buf[0] = '\0';
{0}
return 0;
}}
"""
.. raw:: html
Show Value
.. code-block:: python
"""
extern "C" int call({}) {{
{}
return 0;
}}
"""
.. raw:: html
Show Value
.. code-block:: python
"""
cudaStreamAttrValue stream_attribute;
size_t init_persisting_l2_cache_size;
cudaDeviceGetLimit(&init_persisting_l2_cache_size, cudaLimitPersistingL2CacheSize);
"""
.. raw:: html
Show Value
.. code-block:: python
"""
stream_attribute.accessPolicyWindow.hitRatio = {1};
stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, {2});
stream_attribute.accessPolicyWindow.base_ptr = (void*)({0});
stream_attribute.accessPolicyWindow.num_bytes = {2};
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
"""
.. raw:: html
Show Value
.. code-block:: python
"""
stream_attribute.accessPolicyWindow.num_bytes = 0;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
cudaCtxResetPersistingL2Cache();
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, init_persisting_l2_cache_size);
"""
.. raw:: html
Show Value
.. code-block:: python
"""
CUtensorMap {0};
CUtensorMapDataType {0}_type= (CUtensorMapDataType){1};
cuuint32_t {0}_tensorRank= {2};
void *{0}_globalAddress= {3};
cuuint64_t {0}_globalDim[{2}]= {{{4}}};
cuuint64_t {0}_globalStride[{2}]= {{{5}}};
cuuint32_t {0}_boxDim[{2}]= {{{6}}};
cuuint32_t {0}_elementStrides[{2}]= {{{7}}};
CUtensorMapInterleave {0}_interleave= (CUtensorMapInterleave){8};
CUtensorMapSwizzle {0}_swizzle= (CUtensorMapSwizzle){9};
CUtensorMapL2promotion {0}_l2Promotion= (CUtensorMapL2promotion){10};
CUtensorMapFloatOOBfill {0}_oobFill= (CUtensorMapFloatOOBfill){11};
CUresult {0}_result = CUTLASS_CUDA_DRIVER_WRAPPER_CALL(cuTensorMapEncodeTiled)(
&{0}, {0}_type, {0}_tensorRank, {0}_globalAddress, {0}_globalDim, {0}_globalStride + 1, {0}_boxDim, {0}_elementStrides, {0}_interleave, {0}_swizzle, {0}_l2Promotion, {0}_oobFill);
if ({0}_result != CUDA_SUCCESS) {{
std::stringstream ss;
ss << "Error: Failed to initialize the TMA descriptor {0}";
snprintf(error_buf, ERROR_BUF_SIZE, "%s", ss.str().c_str());
return -1;
}}
"""
.. raw:: html
Show Value
.. code-block:: python
"""
CUtensorMap {0};
CUtensorMapDataType {0}_type= (CUtensorMapDataType){1};
cuuint32_t {0}_tensorRank= {2};
void *{0}_globalAddress= {3};
cuuint64_t {0}_globalDim[{2}]= {{{4}}};
cuuint64_t {0}_globalStride[{2}]= {{{5}}};
cuuint32_t {0}_elementStrides[{2}]= {{{6}}};
int {0}_lowerCorner[{2} - 2]= {{{7}}};
int {0}_upperCorner[{2} - 2]= {{{8}}};
cuuint32_t {0}_channelsPerPixel= {9};
cuuint32_t {0}_pixelsPerColumn= {10};
CUtensorMapInterleave {0}_interleave= (CUtensorMapInterleave){11};
CUtensorMapSwizzle {0}_swizzle= (CUtensorMapSwizzle){12};
CUtensorMapL2promotion {0}_l2Promotion= (CUtensorMapL2promotion){13};
CUtensorMapFloatOOBfill {0}_oobFill= (CUtensorMapFloatOOBfill){14};
CUresult {0}_result = CUTLASS_CUDA_DRIVER_WRAPPER_CALL(cuTensorMapEncodeIm2col)(
&{0}, {0}_type, {0}_tensorRank, {0}_globalAddress, {0}_globalDim, {0}_globalStride + 1,
{0}_lowerCorner, {0}_upperCorner, {0}_channelsPerPixel, {0}_pixelsPerColumn, {0}_elementStrides, {0}_interleave, {0}_swizzle, {0}_l2Promotion, {0}_oobFill);
if ({0}_result != CUDA_SUCCESS) {{
std::stringstream ss;
ss << "Error: Failed to initialize the TMA descriptor {0}";
snprintf(error_buf, ERROR_BUF_SIZE, "%s", ss.str().c_str());
return -1;
}}
"""
.. raw:: html
Show Value
.. code-block:: python
"""
#define ERROR_BUF_SIZE 1024
static char error_buf[ERROR_BUF_SIZE];
extern "C" const char* get_last_error() {
return error_buf;
}
extern "C" int init() {
error_buf[0] = '\0';
return 0;
}
"""
.. raw:: html