[clang] [clang][CodeGen][SPIR-V] Fix incorrect SYCL usage, implement missing interface (PR #109415)
Alex Voicu via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 20 05:40:57 PDT 2024
https://github.com/AlexVlx updated https://github.com/llvm/llvm-project/pull/109415
>From 75ca598c7e8a583545f50ee2c526556df261cc7f Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Fri, 20 Sep 2024 13:25:49 +0100
Subject: [PATCH 1/2] Implement `getGlobalVarAddressSpace` for SPIR-V; stop
using SYCL for tests.
---
clang/lib/Basic/Targets/SPIR.h | 5 ++++
clang/lib/CodeGen/Targets/SPIR.cpp | 23 ++++++++++++++++
.../CodeGenCXX/dynamic-cast-address-space.cpp | 23 ++++++++--------
.../test/CodeGenCXX/spirv-amdgcn-float16.cpp | 27 ++++++++++---------
.../template-param-objects-address-space.cpp | 4 +--
...w-expression-typeinfo-in-address-space.cpp | 2 +-
.../try-catch-with-address-space.cpp | 6 ++---
.../typeid-cxx11-with-address-space.cpp | 2 +-
.../CodeGenCXX/typeid-with-address-space.cpp | 4 +--
.../typeinfo-with-address-space.cpp | 14 +++++-----
.../vtable-assume-load-address-space.cpp | 22 +++++++--------
...e-pointer-initialization-address-space.cpp | 2 +-
clang/test/CodeGenCXX/vtt-address-space.cpp | 2 +-
.../CodeGenOpenCL/builtins-amdgcn-gfx11.cl | 4 +++
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 25 +++++++++++------
15 files changed, 105 insertions(+), 60 deletions(-)
diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h
index 37cf9d7921bac5..4dbc64d27ba909 100644
--- a/clang/lib/Basic/Targets/SPIR.h
+++ b/clang/lib/Basic/Targets/SPIR.h
@@ -380,6 +380,7 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
PointerWidth = PointerAlign = 64;
SizeType = TargetInfo::UnsignedLong;
PtrDiffType = IntPtrType = TargetInfo::SignedLong;
+ AddrSpaceMap = &SPIRDefIsGenMap;
resetDataLayout("e-i64:64-v16:16-v24:32-v32:32-v48:64-"
"v96:128-v192:256-v256:256-v512:512-v1024:1024-G1-P4-A0");
@@ -412,6 +413,10 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
void setAuxTarget(const TargetInfo *Aux) override;
+ void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override {
+ TargetInfo::adjust(Diags, Opts);
+ }
+
bool hasInt128Type() const override { return TargetInfo::hasInt128Type(); }
};
diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index cc52925e2e523f..975b48afa03fcd 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -58,6 +58,8 @@ class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo {
SPIRVTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
: CommonSPIRTargetCodeGenInfo(std::make_unique<SPIRVABIInfo>(CGT)) {}
void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
+ LangAS getGlobalVarAddressSpace(CodeGenModule &CGM,
+ const VarDecl *D) const override;
};
} // End anonymous namespace.
@@ -188,6 +190,27 @@ void SPIRVTargetCodeGenInfo::setCUDAKernelCallingConvention(
}
}
+LangAS SPIRVTargetCodeGenInfo::getGlobalVarAddressSpace(
+ CodeGenModule &CGM, const VarDecl *D) const {
+ assert(!CGM.getLangOpts().OpenCL &&
+ !(CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) &&
+ "Address space agnostic languages only");
+ // If we're here it means that we're using the SPIRDefIsGen ASMap, hence for
+ // the global AS we can rely on either cuda_device or sycl_global to be
+ // correct; however, since this is not a CUDA Device context, we use
+ // sycl_global to prevent confusion with the assertion.
+ LangAS DefaultGlobalAS = getLangASFromTargetAS(
+ CGM.getContext().getTargetAddressSpace(LangAS::sycl_global));
+ if (!D)
+ return DefaultGlobalAS;
+
+ LangAS AddrSpace = D->getType().getAddressSpace();
+ if (AddrSpace != LangAS::Default)
+ return AddrSpace;
+
+ return DefaultGlobalAS;
+}
+
/// Construct a SPIR-V target extension type for the given OpenCL image type.
static llvm::Type *getSPIRVImageType(llvm::LLVMContext &Ctx, StringRef BaseType,
StringRef OpenCLName,
diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
index 3d5e32516c7af2..b967701ca1fa96 100644
--- a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
+++ b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
@@ -1,6 +1,6 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --no-generate-body-for-unused-prefixes --version 4
// RUN: %clang_cc1 -I%S %s -triple amdgcn-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s
-// RUN: %clang_cc1 -I%S %s -triple spirv64-unknown-unknown -fsycl-is-device -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 -I%S %s -triple spirv64-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
struct A { virtual void f(); };
struct B : A { };
@@ -15,7 +15,7 @@ B fail;
// CHECK: @_ZTI1B = linkonce_odr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) getelementptr inbounds (ptr addrspace(1), ptr addrspace(1) @_ZTVN10__cxxabiv120__si_class_type_infoE, i64 2), ptr addrspace(1) @_ZTS1B, ptr addrspace(1) @_ZTI1A }, comdat, align 8
// CHECK: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500
//.
-// WITH-NONZERO-DEFAULT-AS: @_ZTV1B = linkonce_odr unnamed_addr addrspace(1) constant { [3 x ptr addrspace(1)] } { [3 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @_ZTI1B, ptr addrspace(1) addrspacecast (ptr @_ZN1A1fEv to ptr addrspace(1))] }, comdat, align 8
+// WITH-NONZERO-DEFAULT-AS: @_ZTV1B = linkonce_odr unnamed_addr addrspace(1) constant { [3 x ptr addrspace(1)] } { [3 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @_ZTI1B, ptr addrspace(1) addrspacecast (ptr addrspace(4) @_ZN1A1fEv to ptr addrspace(1))] }, comdat, align 8
// WITH-NONZERO-DEFAULT-AS: @fail = addrspace(1) global { ptr addrspace(1) } { ptr addrspace(1) getelementptr inbounds inrange(-16, 8) ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1B, i32 0, i32 0, i32 2) }, align 8
// WITH-NONZERO-DEFAULT-AS: @_ZTI1A = external addrspace(1) constant ptr addrspace(1)
// WITH-NONZERO-DEFAULT-AS: @_ZTVN10__cxxabiv120__si_class_type_infoE = external addrspace(1) global [0 x ptr addrspace(1)]
@@ -60,7 +60,7 @@ B fail;
// CHECK-NEXT: ret ptr addrspacecast (ptr addrspace(1) @fail to ptr)
//
// WITH-NONZERO-DEFAULT-AS-LABEL: define spir_func noundef align 8 dereferenceable(8) ptr addrspace(4) @_Z1fP1A(
-// WITH-NONZERO-DEFAULT-AS-SAME: ptr addrspace(4) noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] personality ptr @__gxx_personality_v0 {
+// WITH-NONZERO-DEFAULT-AS-SAME: ptr addrspace(4) noundef [[A:%.*]]) addrspace(4) #[[ATTR0:[0-9]+]] personality ptr addrspace(4) @__gxx_personality_v0 {
// WITH-NONZERO-DEFAULT-AS-NEXT: entry:
// WITH-NONZERO-DEFAULT-AS-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(4), align 8
// WITH-NONZERO-DEFAULT-AS-NEXT: [[A_ADDR:%.*]] = alloca ptr addrspace(4), align 8
@@ -70,11 +70,11 @@ B fail;
// WITH-NONZERO-DEFAULT-AS-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr [[A_ADDR]] to ptr addrspace(4)
// WITH-NONZERO-DEFAULT-AS-NEXT: store ptr addrspace(4) [[A]], ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// WITH-NONZERO-DEFAULT-AS-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
-// WITH-NONZERO-DEFAULT-AS-NEXT: [[TMP1:%.*]] = call spir_func ptr addrspace(4) @__dynamic_cast(ptr addrspace(4) [[TMP0]], ptr addrspace(1) @_ZTI1A, ptr addrspace(1) @_ZTI1B, i64 0) #[[ATTR3:[0-9]+]]
+// WITH-NONZERO-DEFAULT-AS-NEXT: [[TMP1:%.*]] = call spir_func addrspace(4) ptr addrspace(4) @__dynamic_cast(ptr addrspace(4) [[TMP0]], ptr addrspace(1) @_ZTI1A, ptr addrspace(1) @_ZTI1B, i64 0) #[[ATTR3:[0-9]+]]
// WITH-NONZERO-DEFAULT-AS-NEXT: [[TMP2:%.*]] = icmp eq ptr addrspace(4) [[TMP1]], null
// WITH-NONZERO-DEFAULT-AS-NEXT: br i1 [[TMP2]], label [[DYNAMIC_CAST_BAD_CAST:%.*]], label [[DYNAMIC_CAST_END:%.*]]
// WITH-NONZERO-DEFAULT-AS: dynamic_cast.bad_cast:
-// WITH-NONZERO-DEFAULT-AS-NEXT: invoke spir_func void @__cxa_bad_cast() #[[ATTR4:[0-9]+]]
+// WITH-NONZERO-DEFAULT-AS-NEXT: invoke spir_func addrspace(4) void @__cxa_bad_cast() #[[ATTR4:[0-9]+]]
// WITH-NONZERO-DEFAULT-AS-NEXT: to label [[INVOKE_CONT:%.*]] unwind label [[LPAD:%.*]]
// WITH-NONZERO-DEFAULT-AS: invoke.cont:
// WITH-NONZERO-DEFAULT-AS-NEXT: unreachable
@@ -90,8 +90,8 @@ B fail;
// WITH-NONZERO-DEFAULT-AS-NEXT: br label [[CATCH:%.*]]
// WITH-NONZERO-DEFAULT-AS: catch:
// WITH-NONZERO-DEFAULT-AS-NEXT: [[EXN:%.*]] = load ptr addrspace(4), ptr [[EXN_SLOT]], align 8
-// WITH-NONZERO-DEFAULT-AS-NEXT: [[TMP6:%.*]] = call spir_func ptr addrspace(4) @__cxa_begin_catch(ptr addrspace(4) [[EXN]]) #[[ATTR3]]
-// WITH-NONZERO-DEFAULT-AS-NEXT: call spir_func void @__cxa_end_catch()
+// WITH-NONZERO-DEFAULT-AS-NEXT: [[TMP6:%.*]] = call spir_func addrspace(4) ptr addrspace(4) @__cxa_begin_catch(ptr addrspace(4) [[EXN]]) #[[ATTR3]]
+// WITH-NONZERO-DEFAULT-AS-NEXT: call spir_func addrspace(4) void @__cxa_end_catch()
// WITH-NONZERO-DEFAULT-AS-NEXT: br label [[TRY_CONT]]
// WITH-NONZERO-DEFAULT-AS: try.cont:
// WITH-NONZERO-DEFAULT-AS-NEXT: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @fail to ptr addrspace(4))
@@ -112,9 +112,9 @@ const B& f(A *a) {
// CHECK: attributes #[[ATTR3]] = { nounwind }
// CHECK: attributes #[[ATTR4]] = { noreturn }
//.
-// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32,+wavefrontsize64" }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) }
-// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gws,+image-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32,+wavefrontsize64" }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn }
//.
@@ -122,6 +122,7 @@ const B& f(A *a) {
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
// CHECK: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
//.
-// WITH-NONZERO-DEFAULT-AS: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
-// WITH-NONZERO-DEFAULT-AS: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
+// WITH-NONZERO-DEFAULT-AS: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500}
+// WITH-NONZERO-DEFAULT-AS: [[META1:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}
+// WITH-NONZERO-DEFAULT-AS: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
//.
diff --git a/clang/test/CodeGenCXX/spirv-amdgcn-float16.cpp b/clang/test/CodeGenCXX/spirv-amdgcn-float16.cpp
index 2487e0fcd4343a..223e408a378920 100644
--- a/clang/test/CodeGenCXX/spirv-amdgcn-float16.cpp
+++ b/clang/test/CodeGenCXX/spirv-amdgcn-float16.cpp
@@ -7,22 +7,25 @@
// CHECK-NEXT: [[X:%.*]] = alloca half, align 2
// CHECK-NEXT: [[Y:%.*]] = alloca half, align 2
// CHECK-NEXT: [[Z:%.*]] = alloca half, align 2
-// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr [[X]], align 2
-// CHECK-NEXT: [[TMP1:%.*]] = load half, ptr [[Y]], align 2
+// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)
+// CHECK-NEXT: [[Y_ASCAST:%.*]] = addrspacecast ptr [[Y]] to ptr addrspace(4)
+// CHECK-NEXT: [[Z_ASCAST:%.*]] = addrspacecast ptr [[Z]] to ptr addrspace(4)
+// CHECK-NEXT: [[TMP0:%.*]] = load half, ptr addrspace(4) [[X_ASCAST]], align 2
+// CHECK-NEXT: [[TMP1:%.*]] = load half, ptr addrspace(4) [[Y_ASCAST]], align 2
// CHECK-NEXT: [[ADD:%.*]] = fadd half [[TMP0]], [[TMP1]]
-// CHECK-NEXT: store half [[ADD]], ptr [[Z]], align 2
-// CHECK-NEXT: [[TMP2:%.*]] = load half, ptr [[X]], align 2
-// CHECK-NEXT: [[TMP3:%.*]] = load half, ptr [[Y]], align 2
+// CHECK-NEXT: store half [[ADD]], ptr addrspace(4) [[Z_ASCAST]], align 2
+// CHECK-NEXT: [[TMP2:%.*]] = load half, ptr addrspace(4) [[X_ASCAST]], align 2
+// CHECK-NEXT: [[TMP3:%.*]] = load half, ptr addrspace(4) [[Y_ASCAST]], align 2
// CHECK-NEXT: [[SUB:%.*]] = fsub half [[TMP2]], [[TMP3]]
-// CHECK-NEXT: store half [[SUB]], ptr [[Z]], align 2
-// CHECK-NEXT: [[TMP4:%.*]] = load half, ptr [[X]], align 2
-// CHECK-NEXT: [[TMP5:%.*]] = load half, ptr [[Y]], align 2
+// CHECK-NEXT: store half [[SUB]], ptr addrspace(4) [[Z_ASCAST]], align 2
+// CHECK-NEXT: [[TMP4:%.*]] = load half, ptr addrspace(4) [[X_ASCAST]], align 2
+// CHECK-NEXT: [[TMP5:%.*]] = load half, ptr addrspace(4) [[Y_ASCAST]], align 2
// CHECK-NEXT: [[MUL:%.*]] = fmul half [[TMP4]], [[TMP5]]
-// CHECK-NEXT: store half [[MUL]], ptr [[Z]], align 2
-// CHECK-NEXT: [[TMP6:%.*]] = load half, ptr [[X]], align 2
-// CHECK-NEXT: [[TMP7:%.*]] = load half, ptr [[Y]], align 2
+// CHECK-NEXT: store half [[MUL]], ptr addrspace(4) [[Z_ASCAST]], align 2
+// CHECK-NEXT: [[TMP6:%.*]] = load half, ptr addrspace(4) [[X_ASCAST]], align 2
+// CHECK-NEXT: [[TMP7:%.*]] = load half, ptr addrspace(4) [[Y_ASCAST]], align 2
// CHECK-NEXT: [[DIV:%.*]] = fdiv half [[TMP6]], [[TMP7]]
-// CHECK-NEXT: store half [[DIV]], ptr [[Z]], align 2
+// CHECK-NEXT: store half [[DIV]], ptr addrspace(4) [[Z_ASCAST]], align 2
// CHECK-NEXT: ret void
//
void f() {
diff --git a/clang/test/CodeGenCXX/template-param-objects-address-space.cpp b/clang/test/CodeGenCXX/template-param-objects-address-space.cpp
index b3733decdb550f..1f7c160bb86e91 100644
--- a/clang/test/CodeGenCXX/template-param-objects-address-space.cpp
+++ b/clang/test/CodeGenCXX/template-param-objects-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -std=c++20 %s -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 -triple spirv64-unknown-unknown -fsycl-is-device -std=c++20 %s -emit-llvm -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -std=c++20 %s -emit-llvm -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
struct S { char buf[32]; };
template<S s> constexpr const char *begin() { return s.buf; }
@@ -37,6 +37,6 @@ const void *s = observable_addr<S{"hello world"}>();
// CHECK: define linkonce_odr noundef ptr @_Z15observable_addrIXtl1StlA32_cLc104ELc101ELc108ELc108ELc111ELc32ELc119ELc111ELc114ELc108ELc100EEEEEPKvv()
// WITH-NONZERO-DEFAULT-AS: define linkonce_odr {{.*}} noundef ptr addrspace(4) @_Z15observable_addrIXtl1StlA32_cLc104ELc101ELc108ELc108ELc111ELc32ELc119ELc111ELc114ELc108ELc100EEEEEPKvv()
// CHECK: %call = call noundef ptr @_Z6calleePK1S(ptr noundef addrspacecast (ptr addrspace(1) [[HELLO]] to ptr))
-// WITH-NONZERO-DEFAULT-AS: %call = call {{.*}} noundef ptr addrspace(4) @_Z6calleePK1S(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) [[HELLO]] to ptr addrspace(4)))
+// WITH-NONZERO-DEFAULT-AS: %call = call {{.*}} noundef{{.*}} ptr addrspace(4) @_Z6calleePK1S(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) [[HELLO]] to ptr addrspace(4)))
// CHECK: declare noundef ptr @_Z6calleePK1S(ptr noundef)
// WITH-NONZERO-DEFAULT-AS: declare {{.*}} noundef ptr addrspace(4) @_Z6calleePK1S(ptr addrspace(4) noundef)
diff --git a/clang/test/CodeGenCXX/throw-expression-typeinfo-in-address-space.cpp b/clang/test/CodeGenCXX/throw-expression-typeinfo-in-address-space.cpp
index 3acbdd8fd97ee0..0fb553de4485af 100644
--- a/clang/test/CodeGenCXX/throw-expression-typeinfo-in-address-space.cpp
+++ b/clang/test/CodeGenCXX/throw-expression-typeinfo-in-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -std=c++11 -o - | FileCheck %s
-// RUN: %clang_cc1 %s -triple spirv64-unknown-unknown -fsycl-is-device -emit-llvm -fcxx-exceptions -fexceptions -std=c++11 -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 %s -triple spirv64-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -std=c++11 -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
struct X {
~X();
diff --git a/clang/test/CodeGenCXX/try-catch-with-address-space.cpp b/clang/test/CodeGenCXX/try-catch-with-address-space.cpp
index 412ac6c2872587..55c76ed7f344f9 100644
--- a/clang/test/CodeGenCXX/try-catch-with-address-space.cpp
+++ b/clang/test/CodeGenCXX/try-catch-with-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -o - -fcxx-exceptions -fexceptions | FileCheck %s
-// RUN: %clang_cc1 %s -triple=spirv64-unknown-unknown -fsycl-is-device -emit-llvm -o - -fcxx-exceptions -fexceptions | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 %s -triple=spirv64-amd-amdhsa -emit-llvm -o - -fcxx-exceptions -fexceptions | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
struct X { };
@@ -12,7 +12,7 @@ void f() {
} catch (const X x) {
// CHECK: catch ptr addrspace(1) @_ZTI1X
// CHECK: call i32 @llvm.eh.typeid.for.p0(ptr addrspacecast (ptr addrspace(1) @_ZTI1X to ptr))
- // WITH-NONZERO-DEFAULT-AS: call i32 @llvm.eh.typeid.for.p4(ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZTI1X to ptr addrspace(4)))
+ // WITH-NONZERO-DEFAULT-AS: call{{.*}} i32 @llvm.eh.typeid.for.p4(ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZTI1X to ptr addrspace(4)))
}
}
@@ -23,6 +23,6 @@ void h() {
} catch (char const(&)[4]) {
// CHECK: catch ptr addrspace(1) @_ZTIA4_c
// CHECK: call i32 @llvm.eh.typeid.for.p0(ptr addrspacecast (ptr addrspace(1) @_ZTIA4_c to ptr))
- // WITH-NONZERO-DEFAULT-AS: call i32 @llvm.eh.typeid.for.p4(ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZTIA4_c to ptr addrspace(4)))
+ // WITH-NONZERO-DEFAULT-AS: call{{.*}} i32 @llvm.eh.typeid.for.p4(ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZTIA4_c to ptr addrspace(4)))
}
}
diff --git a/clang/test/CodeGenCXX/typeid-cxx11-with-address-space.cpp b/clang/test/CodeGenCXX/typeid-cxx11-with-address-space.cpp
index f6dc38ec9f2925..782e59c887bd36 100644
--- a/clang/test/CodeGenCXX/typeid-cxx11-with-address-space.cpp
+++ b/clang/test/CodeGenCXX/typeid-cxx11-with-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -I%S %s -triple amdgcn-amd-amdhsa -emit-llvm -std=c++11 -o - | FileCheck %s
-// RUN: %clang_cc1 -I%S %s -triple spirv64-unknown-unknown -fsycl-is-device -emit-llvm -std=c++11 -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 -I%S %s -triple spirv64-amd-amdhsa -emit-llvm -std=c++11 -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
#include <typeinfo>
namespace Test1 {
diff --git a/clang/test/CodeGenCXX/typeid-with-address-space.cpp b/clang/test/CodeGenCXX/typeid-with-address-space.cpp
index 98af17f4fc8888..20699dc170acea 100644
--- a/clang/test/CodeGenCXX/typeid-with-address-space.cpp
+++ b/clang/test/CodeGenCXX/typeid-with-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -I%S %s -triple amdgcn-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s
-// RUN: %clang_cc1 -I%S %s -triple spirv64-unknown-unknown -fsycl-is-device -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 -I%S %s -triple spirv64-amd-amdhsa -emit-llvm -fcxx-exceptions -fexceptions -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
#include <typeinfo>
namespace Test1 {
@@ -39,7 +39,7 @@ const std::type_info &A10_c_ti = typeid(char const[10]);
// CHECK-LABEL: define{{.*}} ptr @_ZN5Test11fEv
// CHECK-SAME: personality ptr @__gxx_personality_v0
// WITH-NONZERO-DEFAULT-AS-LABEL: define{{.*}} ptr addrspace(4) @_ZN5Test11fEv
-// WITH-NONZERO-DEFAULT-AS-SAME: personality ptr @__gxx_personality_v0
+// WITH-NONZERO-DEFAULT-AS-SAME: personality ptr addrspace(4) @__gxx_personality_v0
const char *f() {
try {
// CHECK: br i1
diff --git a/clang/test/CodeGenCXX/typeinfo-with-address-space.cpp b/clang/test/CodeGenCXX/typeinfo-with-address-space.cpp
index 350303cc6e9b34..60eb8f17f91fd1 100644
--- a/clang/test/CodeGenCXX/typeinfo-with-address-space.cpp
+++ b/clang/test/CodeGenCXX/typeinfo-with-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -I%S %s -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck %s -check-prefix=AS
-// RUN: %clang_cc1 -I%S %s -triple spirv64-unknown-unknown -fsycl-is-device -emit-llvm -o - | FileCheck %s -check-prefix=NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 -I%S %s -triple spirv64-amd-amdhsa -emit-llvm -o - | FileCheck %s -check-prefix=NONZERO-DEFAULT-AS
// RUN: %clang_cc1 -I%S %s -triple x86_64-linux-gnu -emit-llvm -o - | FileCheck %s -check-prefix=NO-AS
#include <typeinfo>
@@ -26,30 +26,30 @@ class B : A {
unsigned long Fn(B& b) {
// AS: %call = call noundef zeroext i1 @_ZNKSt9type_infoeqERKS_(ptr {{.*}} addrspacecast (ptr addrspace(1) @_ZTISt9type_info to ptr), ptr {{.*}} %2)
-// NONZERO-DEFAULT-AS: %call = call{{.*}} noundef zeroext i1 @_ZNKSt9type_infoeqERKS_(ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTISt9type_info to ptr addrspace(4)), ptr addrspace(4) {{.*}} %2)
+// NONZERO-DEFAULT-AS: %call = call{{.*}} noundef zeroext{{.*}} i1 @_ZNKSt9type_infoeqERKS_(ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTISt9type_info to ptr addrspace(4)), ptr addrspace(4) {{.*}} %2)
// NO-AS: %call = call noundef zeroext i1 @_ZNKSt9type_infoeqERKS_(ptr {{.*}} @_ZTISt9type_info, ptr {{.*}} %2)
if (typeid(std::type_info) == typeid(b))
return 42;
// AS: %call2 = call noundef zeroext i1 @_ZNKSt9type_infoneERKS_(ptr {{.*}} addrspacecast (ptr addrspace(1) @_ZTIi to ptr), ptr {{.*}} %5)
-// NONZERO-DEFAULT-AS: %call2 = call{{.*}} noundef zeroext i1 @_ZNKSt9type_infoneERKS_(ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTIi to ptr addrspace(4)), ptr addrspace(4) {{.*}} %5)
+// NONZERO-DEFAULT-AS: %call2 = call{{.*}} noundef zeroext{{.*}} i1 @_ZNKSt9type_infoneERKS_(ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTIi to ptr addrspace(4)), ptr addrspace(4) {{.*}} %5)
// NO-AS: %call2 = call noundef zeroext i1 @_ZNKSt9type_infoneERKS_(ptr {{.*}} @_ZTIi, ptr {{.*}} %5)
if (typeid(int) != typeid(b))
return 1712;
// AS: %call5 = call noundef ptr @_ZNKSt9type_info4nameEv(ptr {{.*}} addrspacecast (ptr addrspace(1) @_ZTI1A to ptr))
-// NONZERO-DEFAULT-AS: %call5 = call{{.*}} noundef ptr addrspace(4) @_ZNKSt9type_info4nameEv(ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTI1A to ptr addrspace(4)))
+// NONZERO-DEFAULT-AS: %call5 = call{{.*}} noundef{{.*}} ptr addrspace(4) @_ZNKSt9type_info4nameEv(ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTI1A to ptr addrspace(4)))
// NO-AS: %call5 = call noundef ptr @_ZNKSt9type_info4nameEv(ptr {{.*}} @_ZTI1A)
// AS: %call7 = call noundef ptr @_ZNKSt9type_info4nameEv(ptr {{.*}} %8)
-// NONZERO-DEFAULT-AS: %call7 = call{{.*}} noundef ptr addrspace(4) @_ZNKSt9type_info4nameEv(ptr addrspace(4) {{.*}} %8)
+// NONZERO-DEFAULT-AS: %call7 = call{{.*}} noundef{{.*}} ptr addrspace(4) @_ZNKSt9type_info4nameEv(ptr addrspace(4) {{.*}} %8)
// NO-AS: %call7 = call noundef ptr @_ZNKSt9type_info4nameEv(ptr {{.*}} %8)
if (typeid(A).name() == typeid(b).name())
return 0;
// AS: %call11 = call noundef zeroext i1 @_ZNKSt9type_info6beforeERKS_(ptr {{.*}} %11, ptr {{.*}} addrspacecast (ptr addrspace(1) @_ZTIf to ptr))
-// NONZERO-DEFAULT-AS: %call11 = call{{.*}} noundef zeroext i1 @_ZNKSt9type_info6beforeERKS_(ptr addrspace(4) {{.*}} %11, ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTIf to ptr addrspace(4)))
+// NONZERO-DEFAULT-AS: %call11 = call{{.*}} noundef zeroext{{.*}} i1 @_ZNKSt9type_info6beforeERKS_(ptr addrspace(4) {{.*}} %11, ptr addrspace(4) {{.*}} addrspacecast (ptr addrspace(1) @_ZTIf to ptr addrspace(4)))
// NO-AS: %call11 = call noundef zeroext i1 @_ZNKSt9type_info6beforeERKS_(ptr {{.*}} %11, ptr {{.*}} @_ZTIf)
if (typeid(b).before(typeid(float)))
return 1;
// AS: %call15 = call noundef i64 @_ZNKSt9type_info9hash_codeEv(ptr {{.*}} %14)
-// NONZERO-DEFAULT-AS: %call15 = call{{.*}} noundef i64 @_ZNKSt9type_info9hash_codeEv(ptr addrspace(4) {{.*}} %14)
+// NONZERO-DEFAULT-AS: %call15 = call{{.*}} noundef{{.*}} i64 @_ZNKSt9type_info9hash_codeEv(ptr addrspace(4) {{.*}} %14)
// NO-AS: %call15 = call noundef i64 @_ZNKSt9type_info9hash_codeEv(ptr {{.*}} %14)
return typeid(b).hash_code();
}
diff --git a/clang/test/CodeGenCXX/vtable-assume-load-address-space.cpp b/clang/test/CodeGenCXX/vtable-assume-load-address-space.cpp
index ecafa99d8be00f..e8d3a8407cda73 100644
--- a/clang/test/CodeGenCXX/vtable-assume-load-address-space.cpp
+++ b/clang/test/CodeGenCXX/vtable-assume-load-address-space.cpp
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o %t.ll -O1 -disable-llvm-passes -fms-extensions -fstrict-vtable-pointers
// RUN: %clang_cc1 %s -triple i686-pc-win32 -emit-llvm -o %t.ms.ll -O1 -disable-llvm-passes -fms-extensions -fstrict-vtable-pointers
-// RUN: %clang_cc1 %s -triple=spirv64-unknown-unknown -fsycl-is-device -std=c++11 -emit-llvm -o %t.ll -O1 -disable-llvm-passes -fms-extensions -fstrict-vtable-pointers
+// RUN: %clang_cc1 %s -triple=spirv64-amd-amdhsa -std=c++11 -emit-llvm -o %t.ll -O1 -disable-llvm-passes -fms-extensions -fstrict-vtable-pointers
// FIXME: Assume load should not require -fstrict-vtable-pointers
// RUN: FileCheck --check-prefix=CHECK1 --input-file=%t.ll %s
@@ -29,7 +29,7 @@ void g(A *a) { a->foo(); }
// CHECK1: call{{.*}} void @_ZN5test11AC1Ev(ptr {{((addrspace(4)){0,1})}}
// CHECK1: %[[VTABLE:.*]] = load ptr addrspace(1), ptr {{((addrspace(4)){0,1})}}{{.*}}%{{.*}}
// CHECK1: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds inrange(-16, 8) ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test11AE, i32 0, i32 0, i32 2)
-// CHECK1: call void @llvm.assume(i1 %[[CMP]])
+// CHECK1: call{{.*}} void @llvm.assume(i1 %[[CMP]])
// CHECK1-LABEL: {{^}}}
void fooA() {
@@ -41,7 +41,7 @@ void fooA() {
// CHECK1: call{{.*}} void @_ZN5test11BC1Ev(ptr {{[^,]*}} %{{.*}})
// CHECK1: %[[VTABLE:.*]] = load ptr addrspace(1), ptr {{((addrspace(4)){0,1})}}{{.*}}%{{.*}}
// CHECK1: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds inrange(-16, 8) ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test11BE, i32 0, i32 0, i32 2)
-// CHECK1: call void @llvm.assume(i1 %[[CMP]])
+// CHECK1: call{{.*}} void @llvm.assume(i1 %[[CMP]])
// CHECK1-LABEL: {{^}}}
void fooB() {
@@ -75,12 +75,12 @@ void h(B *b) { b->bar(); }
// CHECK2: call{{.*}} void @_ZN5test21CC1Ev(ptr
// CHECK2: %[[VTABLE:.*]] = load ptr addrspace(1), ptr {{.*}}
// CHECK2: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds inrange(-16, 8) ({ [3 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test21CE, i32 0, i32 0, i32 2)
-// CHECK2: call void @llvm.assume(i1 %[[CMP]])
+// CHECK2: call{{.*}} void @llvm.assume(i1 %[[CMP]])
// CHECK2: %[[ADD_PTR:.*]] = getelementptr inbounds i8, ptr {{((addrspace(4)){0,1})}}{{.*}}%{{.*}}, i64 8
// CHECK2: %[[VTABLE2:.*]] = load ptr addrspace(1), ptr {{((addrspace(4)){0,1})}}{{.*}}%[[ADD_PTR]]
// CHECK2: %[[CMP2:.*]] = icmp eq ptr addrspace(1) %[[VTABLE2]], getelementptr inbounds inrange(-16, 8) ({ [3 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test21CE, i32 0, i32 1, i32 2)
-// CHECK2: call void @llvm.assume(i1 %[[CMP2]])
+// CHECK2: call{{.*}} void @llvm.assume(i1 %[[CMP2]])
// CHECK2: call{{.*}} void @_ZN5test21gEPNS_1AE(
// CHECK2-LABEL: {{^}}}
@@ -111,7 +111,7 @@ void g(B *a) { a->foo(); }
// CHECK3-LABEL: define{{.*}} void @_ZN5test34testEv()
// CHECK3: call{{.*}} void @_ZN5test31CC1Ev(ptr
// CHECK3: %[[CMP:.*]] = icmp eq ptr addrspace(1) %{{.*}}, getelementptr inbounds inrange(-24, 8) ({ [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test31CE, i32 0, i32 0, i32 3)
-// CHECK3: call void @llvm.assume(i1 %[[CMP]])
+// CHECK3: call{{.*}} void @llvm.assume(i1 %[[CMP]])
// CHECK3-LABLEL: }
void test() {
C c;
@@ -140,11 +140,11 @@ void g(C *c) { c->foo(); }
// CHECK4: call{{.*}} void @_ZN5test41CC1Ev(ptr
// CHECK4: %[[VTABLE:.*]] = load ptr addrspace(1), ptr {{((addrspace(4)){0,1})}}{{.*}}%{{.*}}
// CHECK4: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds inrange(-32, 8) ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test41CE, i32 0, i32 0, i32 4)
-// CHECK4: call void @llvm.assume(i1 %[[CMP]]
+// CHECK4: call{{.*}} void @llvm.assume(i1 %[[CMP]]
// CHECK4: %[[VTABLE2:.*]] = load ptr addrspace(1), ptr {{((addrspace(4)){0,1})}}{{.*}}%{{.*}}
// CHECK4: %[[CMP2:.*]] = icmp eq ptr addrspace(1) %[[VTABLE2]], getelementptr inbounds inrange(-32, 8) ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test41CE, i32 0, i32 0, i32 4)
-// CHECK4: call void @llvm.assume(i1 %[[CMP2]])
+// CHECK4: call{{.*}} void @llvm.assume(i1 %[[CMP2]])
// CHECK4-LABEL: {{^}}}
void test() {
@@ -214,7 +214,7 @@ void A::foo() {}
// CHECK7-LABEL: define{{.*}} void @_ZN5test71gEv()
// CHECK7: call{{.*}} void @_ZN5test71AC1Ev(
-// CHECK7: call void @llvm.assume(
+// CHECK7: call{{.*}} void @llvm.assume(
// CHECK7-LABEL: {{^}}}
void g() {
A *a = new A();
@@ -257,7 +257,7 @@ struct E : A {
};
// CHECK8-LABEL: define{{.*}} void @_ZN5test81bEv()
-// CHECK8: call void @llvm.assume(
+// CHECK8: call{{.*}} void @llvm.assume(
// CHECK8-LABEL: {{^}}}
void b() {
B b;
@@ -285,7 +285,7 @@ void d() {
}
// CHECK8-LABEL: define{{.*}} void @_ZN5test81eEv()
-// CHECK8: call void @llvm.assume(
+// CHECK8: call{{.*}} void @llvm.assume(
// CHECK8-LABEL: {{^}}}
void e() {
E e;
diff --git a/clang/test/CodeGenCXX/vtable-pointer-initialization-address-space.cpp b/clang/test/CodeGenCXX/vtable-pointer-initialization-address-space.cpp
index 876d0845cc5157..8b5b9dd3530648 100644
--- a/clang/test/CodeGenCXX/vtable-pointer-initialization-address-space.cpp
+++ b/clang/test/CodeGenCXX/vtable-pointer-initialization-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 %s -triple=spirv64-unknown-unknown -fsycl-is-device -std=c++11 -emit-llvm -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 %s -triple=spirv64-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
struct Field {
Field();
diff --git a/clang/test/CodeGenCXX/vtt-address-space.cpp b/clang/test/CodeGenCXX/vtt-address-space.cpp
index 4c3d0a534611c7..3409bc7cbb69dd 100644
--- a/clang/test/CodeGenCXX/vtt-address-space.cpp
+++ b/clang/test/CodeGenCXX/vtt-address-space.cpp
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 %s -triple=spirv64-unknown-unknown -fsycl-is-device -std=c++11 -emit-llvm -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
+// RUN: %clang_cc1 %s -triple=spirv64-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s --check-prefix=WITH-NONZERO-DEFAULT-AS
// This is the sample from the C++ Itanium ABI, p2.6.2.
namespace Test {
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
index 138616ccca7182..7f07160bf5e886 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11.cl
@@ -50,6 +50,10 @@ void test_s_wait_event_export_ready() {
// CHECK-LABEL: @test_global_add_f32
// CHECK: = atomicrmw fadd ptr addrspace(1) %addr, float %x syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.ignore.denormal.mode !{{[0-9]+$}}
+#if !defined(__SPIRV__)
void test_global_add_f32(float *rtn, global float *addr, float x) {
+#else
+void test_global_add_f32(float *rtn, __attribute__((address_space(1))) float *addr, float x) {
+#endif
*rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x);
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 6a6d5b1dfed3df..86111f5b4f8584 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -669,23 +669,24 @@ void test_s_getpc(global ulong* out)
}
// CHECK-LABEL: @test_ds_append_lds(
-// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.append.p3(ptr addrspace(3) %ptr, i1 false)
-kernel void test_ds_append_lds(global int* out, local int* ptr) {
+// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.append.p3(ptr addrspace(3) %{{.+}}, i1 false)
#if !defined(__SPIRV__)
- *out = __builtin_amdgcn_ds_append(ptr);
+kernel void test_ds_append_lds(global int* out, local int* ptr) {
#else
- *out = __builtin_amdgcn_ds_append((__attribute__((address_space(3))) int*)(int*)ptr);
+kernel void test_ds_append_lds(__attribute__((address_space(1))) int* out, __attribute__((address_space(3))) int* ptr) {
#endif
+ *out = __builtin_amdgcn_ds_append(ptr);
}
// CHECK-LABEL: @test_ds_consume_lds(
-// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.consume.p3(ptr addrspace(3) %ptr, i1 false)
-kernel void test_ds_consume_lds(global int* out, local int* ptr) {
+// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.consume.p3(ptr addrspace(3) %{{.+}}, i1 false)
+
#if !defined(__SPIRV__)
- *out = __builtin_amdgcn_ds_consume(ptr);
+kernel void test_ds_consume_lds(global int* out, local int* ptr) {
#else
- *out = __builtin_amdgcn_ds_consume((__attribute__((address_space(3))) int*)(int*)ptr);
+kernel void test_ds_consume_lds(__attribute__((address_space(1))) int* out, __attribute__((address_space(3))) int* ptr) {
#endif
+ *out = __builtin_amdgcn_ds_consume(ptr);
}
// CHECK-LABEL: @test_gws_init(
@@ -839,7 +840,11 @@ kernel void test_s_setreg(uint val) {
}
// CHECK-LABEL test_atomic_inc_dec(
+#if !defined(__SPIRV__)
void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
+#else
+void test_atomic_inc_dec(__attribute__((address_space(3))) uint *lptr, __attribute__((address_space(1))) uint *gptr, uint val) {
+#endif
uint res;
// CHECK: atomicrmw uinc_wrap ptr addrspace(3) %lptr, i32 %val syncscope("workgroup") seq_cst, align 4
@@ -855,7 +860,11 @@ void test_atomic_inc_dec(local uint *lptr, global uint *gptr, uint val) {
res = __builtin_amdgcn_atomic_dec32(gptr, val, __ATOMIC_SEQ_CST, "");
// CHECK: atomicrmw volatile udec_wrap ptr addrspace(1) %gptr, i32 %val seq_cst, align 4
+ #if !defined(__SPIRV__)
res = __builtin_amdgcn_atomic_dec32((volatile global uint*)gptr, val, __ATOMIC_SEQ_CST, "");
+ #else
+ res = __builtin_amdgcn_atomic_dec32((volatile __attribute__((address_space(1))) uint*)gptr, val, __ATOMIC_SEQ_CST, "");
+ #endif
}
// CHECK-LABEL test_wavefrontsize(
>From 6b8c55f063cfe297e134980143852ecd3b0e78b9 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Fri, 20 Sep 2024 13:40:43 +0100
Subject: [PATCH 2/2] Fix formatting.
---
clang/lib/CodeGen/Targets/SPIR.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index 975b48afa03fcd..ae87506e535934 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -190,8 +190,9 @@ void SPIRVTargetCodeGenInfo::setCUDAKernelCallingConvention(
}
}
-LangAS SPIRVTargetCodeGenInfo::getGlobalVarAddressSpace(
- CodeGenModule &CGM, const VarDecl *D) const {
+LangAS
+SPIRVTargetCodeGenInfo::getGlobalVarAddressSpace(CodeGenModule &CGM,
+ const VarDecl *D) const {
assert(!CGM.getLangOpts().OpenCL &&
!(CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) &&
"Address space agnostic languages only");
More information about the cfe-commits
mailing list