[DLIGHT][GPU] Enhance opencl thread limit for schedules

python/tvm/dlight/gpu/general_reduction.py

@@ -40,6 +40,9 @@ def apply(  # pylint: disable=too-many-locals
         if target.kind.name == "cuda":
             len_tx = 256
             unroll_depth = 256
+        elif target.kind.name == "opencl":
+            len_tx = 256
+            unroll_depth = 64
         else:
             len_tx = 64
             unroll_depth = 64

python/tvm/dlight/gpu/rmsnorm.py

@@ -82,6 +82,8 @@ def apply(  # pylint: disable=too-many-locals,missing-docstring
     ) -> tir.Schedule:
         if target.kind.name == "cuda":
             num_tx = 512
+        elif target.kind.name == "opencl":
+            num_tx = 256
         else:
             num_tx = 64
 

python/tvm/dlight/gpu/transpose.py

@@ -57,6 +57,10 @@ def apply(  # pylint: disable=too-many-locals
             len_tx = 16
             len_ty = 8
             unroll_depth = 256
+        elif target.kind.name == "opencl":
+            len_tx = 16
+            len_ty = 8
+            unroll_depth = 64
         else:
             len_tx = 8
             len_ty = 4

python/tvm/dlight/gpu/utils.py

@@ -55,6 +55,8 @@ def suggest_threads_per_block(
         threads = 256
     elif target.kind.name == "metal":
         threads = 256
+    elif target.kind.name == "opencl":
+        threads = 256
     else:
         threads = 64
     results: List[Optional[int]] = []