add a passing test

Author: tenpercent

Jenkinsfile

@@ -573,7 +573,7 @@ def Build_CK(Map conf=[:]){
                                   . ${env.WORKSPACE}/bin/activate
                                   python3 -m pip install pytest build setuptools setuptools_scm
                                   python3 -m pip install .
-                                  python3 -m pytest python/test/test_gen_instances.py
+                                  python3 -m pytest python/test/ck4inductor/test_gen_instances.py
                             """
                     }
                     dir("build"){

pyproject.toml

@@ -7,7 +7,7 @@ name = "rocm-composable-kernel"
 dynamic = ["version"]
 description = "Composable Kernel, performance-critical kernels for machine learning workloads"
 readme = "README.md"
-requires-python = ">=3.8"
+requires-python = ">=3.10"
 license = {file = "LICENSE"}
 classifiers = [
     "Programming Language :: Python :: 3",
@@ -21,8 +21,13 @@ dependencies = []
 "Bug Tracker" = "https://github.com/rocm/composable_kernel/issues"
 
 [tool.setuptools]
+
 packages = [
     "genck", 
+    "genck.ops",
+    "genck.ops.ck_tile",
+    "genck.ops.ck_tile.gemm",
+    "genck.emitters",
     "ck4inductor", 
     "ck4inductor.include", 
     "ck4inductor.library", 
@@ -32,18 +37,26 @@ packages = [
 ]
 
 [tool.setuptools.package-dir]
+
+genck = "python/codegen"
+"genck.ops" = "python/codegen/ops"
+"genck.ops.ck_tile" = "python/codegen/ops/ck_tile"
+"genck.ops.ck_tile.gemm" = "python/codegen/ops/ck_tile/gemm"
+"genck.emitters" = "python/codegen/emitters"
 ck4inductor = "python/ck4inductor"
 "ck4inductor.universal_gemm" = "python/ck4inductor/universal_gemm"
 "ck4inductor.batched_universal_gemm" = "python/ck4inductor/batched_universal_gemm"
 "ck4inductor.grouped_conv_fwd" = "python/ck4inductor/grouped_conv_fwd"
 "ck4inductor.include" = "include"
 "ck4inductor.library" = "library"
 
-"genck" = "python/codegen"
-
 [tool.setuptools.package-data]
 "ck4inductor.include" = ["ck/**/*.hpp"]
-"ck4inductor.library" = ["src/tensor_operation_instance/gpu/gemm_universal/**/*.hpp", "src/tensor_operation_instance/gpu/gemm_universal_batched/**/*.hpp", "include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/**/*.hpp"]
+"ck4inductor.library" = [
+    "src/tensor_operation_instance/gpu/gemm_universal/device_gemm_xdl_universal_*/*.hpp", 
+    "src/tensor_operation_instance/gpu/gemm_universal_batched/**/*.hpp", 
+    "include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/**/*.hpp"
+]
 
 [tool.setuptools.dynamic]
 version = { attr = "setuptools_scm.get_version" }

python/codegen/__init__.py

@@ -0,0 +1,33 @@
+import unittest
+
+class TestCanRenderInstance(unittest.TestCase):
+    def test_can_render_instance(self):
+        from genck.ops.ck_tile.gemm.instance import GEMM
+        from genck.ops.ck_tile.gemm.render import render
+
+        test_instance = GEMM(
+            layout_a="Row",
+            layout_b="Col",
+            layout_c="Row",
+            datatype_a="BF16",
+            datatype_b="BF16",
+            datatype_c="BF16",
+            tile_m=256,
+            tile_n=256,
+            tile_k=64,
+            warp_m=2,
+            warp_n=2,
+            warp_k=1,
+            warp_tile_m=32,
+            warp_tile_n=32,
+            warp_tile_k=16,
+            m_is_padded="false",
+            n_is_padded="false",
+            k_is_padded="false",
+            pipeline="CompV3",
+            scheduler="Intrawave",
+            epilogue="Default",
+        )
+
+        rendered_instance = render(test_instance)
+        self.assertIn("ck_tile_gemm_universal", rendered_instance)