This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 7bea15f162 [WINDOWS] Compiler options for non x86 targets (#17260)
7bea15f162 is described below
commit 7bea15f162ceb3f38809212eec5d711929709620
Author: krishnaraj36 <[email protected]>
AuthorDate: Wed Aug 21 00:53:53 2024 +0530
[WINDOWS] Compiler options for non x86 targets (#17260)
---
python/tvm/contrib/cc.py | 5 ++++-
python/tvm/dlight/gpu/gemv.py | 15 +++++++++++----
2 files changed, 15 insertions(+), 5 deletions(-)
diff --git a/python/tvm/contrib/cc.py b/python/tvm/contrib/cc.py
index 59b57e08ba..110f80db61 100644
--- a/python/tvm/contrib/cc.py
+++ b/python/tvm/contrib/cc.py
@@ -372,8 +372,11 @@ def _linux_compile(
def _windows_compile(output, objects, options, cwd=None, ccache_env=None):
- cmd = ["clang"]
+ compiler = os.getenv("TVM_WIN_CC", default="clang")
+ win_target = os.getenv("TVM_WIN_TARGET", default="x86_64")
+ cmd = [compiler]
cmd += ["-O2"]
+ cmd += ["--target=" + win_target]
if output.endswith(".so") or output.endswith(".dll"):
cmd += ["-shared"]
diff --git a/python/tvm/dlight/gpu/gemv.py b/python/tvm/dlight/gpu/gemv.py
index 2bcb8563a2..cff234140e 100644
--- a/python/tvm/dlight/gpu/gemv.py
+++ b/python/tvm/dlight/gpu/gemv.py
@@ -11,7 +11,7 @@
# Unless required by applicable law or agreed to in writing,
# software distributed under the License is distributed on an
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
+# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
"""A rule for GEMV and DecodeGEMV."""
@@ -478,7 +478,9 @@ class GEMV(GPUScheduleRule):
TS, TR = 8, 64
else:
TS, TR = 1, 64
- elif target.kind.name == "opencl" and "android" in str(target.host):
+ elif target.kind.name == "opencl" and (
+ ("android" in str(target.host)) or ("adreno" in str(target.attrs))
+ ):
TAG_S, TAG_R = "threadIdx.x", "threadIdx.y"
VEC_C = 8
LOAD_V_SHARED = False
@@ -686,7 +688,9 @@ class GEMV(GPUScheduleRule):
DEC_PACK = 8
SCALE_PACK = 4
- if target.kind.name == "opencl" and "android" in str(target.host):
+ if target.kind.name == "opencl" and (
+ ("android" in str(target.host)) or ("adreno" in str(target.attrs))
+ ):
TAG_S, TAG_R = "threadIdx.x", "threadIdx.y"
VEC_C = 8
UNROLL = 8
@@ -756,7 +760,10 @@ class GEMV(GPUScheduleRule):
):
"""Schedule the outer reduction block."""
# NOTE: Only Android is supported so far
- if not (target.kind.name == "opencl" and "android" in
str(target.host)):
+ if not (
+ target.kind.name == "opencl"
+ and (("android" in str(target.host)) or ("adreno" in
str(target.attrs)))
+ ):
return None
batch, s, r, c = sch.get_loops(block)
len_s = get_extent(sch, s)
