From 315a70be59029368688a76d47cb4e55de3c4b917 Mon Sep 17 00:00:00 2001
From: Harald van Dijk <harald.vandijk@codeplay.com>
Date: Wed, 22 Jan 2025 15:53:38 +0000
Subject: [PATCH] [vecz] Handle missing gather/scatter functions.

If we encounter types we do not know how to mangle, we cannot generate
gather or scatter functions. Just use scalar loads and stores for this.
---
 .../vecz/source/transform/packetizer.cpp      |  2 +
 .../vecz/test/lit/llvm/packetize_i48.ll       | 50 +++++++++++++++++++
 2 files changed, 52 insertions(+)
 create mode 100644 modules/compiler/vecz/test/lit/llvm/packetize_i48.ll

diff --git a/modules/compiler/vecz/source/transform/packetizer.cpp b/modules/compiler/vecz/source/transform/packetizer.cpp
index 2299c8a22..2e68e5330 100644
--- a/modules/compiler/vecz/source/transform/packetizer.cpp
+++ b/modules/compiler/vecz/source/transform/packetizer.cpp
@@ -2642,12 +2642,14 @@ ValuePacket Packetizer::Impl::packetizeMemOp(MemOp &op) {
         auto *gather =
             createGather(Ctx, packetVecTy, ptrPacket[i], maskPacket[i], EVL,
                          op.getAlignment(), name);
+        PACK_FAIL_IF(!gather);
         gather->insertBefore(op.getInstr()->getIterator());
         results.push_back(gather);
       } else {
         auto *scatter =
             createScatter(Ctx, dataPacket[i], ptrPacket[i], maskPacket[i], EVL,
                           op.getAlignment(), name);
+        PACK_FAIL_IF(!scatter);
         scatter->insertBefore(op.getInstr()->getIterator());
         results.push_back(scatter);
       }
diff --git a/modules/compiler/vecz/test/lit/llvm/packetize_i48.ll b/modules/compiler/vecz/test/lit/llvm/packetize_i48.ll
new file mode 100644
index 000000000..d86793e43
--- /dev/null
+++ b/modules/compiler/vecz/test/lit/llvm/packetize_i48.ll
@@ -0,0 +1,50 @@
+; Copyright (C) Codeplay Software Limited
+;
+; Licensed under the Apache License, Version 2.0 (the "License") with LLVM
+; Exceptions; you may not use this file except in compliance with the License.
+; You may obtain a copy of the License at
+;
+;     https://github.com/codeplaysoftware/oneapi-construction-kit/blob/main/LICENSE.txt
+;
+; 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 specific language governing permissions and limitations
+; under the License.
+;
+; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+; RUN: veczc -k test -vecz-simd-width=4 -S < %s | FileCheck %s
+
+target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "spir64-unknown-unknown"
+
+declare i64 @__mux_get_local_id(i32)
+
+define spir_kernel void @test(ptr %0, ptr %1) {
+entry:
+  %lid = tail call i64 @__mux_get_local_id(i32 0)
+  %ptr.0 = getelementptr i32, ptr %0, i64 %lid
+  %ptr.1 = getelementptr i32, ptr %1, i64 %lid
+  %val = load i48, ptr %ptr.0
+  store i48 %val, ptr %ptr.1
+  ret void
+}
+
+; CHECK-LABEL: define spir_kernel void @test
+; CHECK: load i48
+; CHECK-NOT: load i48
+; CHECK: store i48
+; CHECK-NOT: store i48
+
+; CHECK-LABEL: define spir_kernel void @__vecz_v4_test
+; CHECK: load i48
+; CHECK: load i48
+; CHECK: load i48
+; CHECK: load i48
+; CHECK-NOT: load i48
+; CHECK: store i48
+; CHECK: store i48
+; CHECK: store i48
+; CHECK: store i48
+; CHECK-NOT: store i48