[llvm] bcad161 - [Clang][SPIR-V] Emit target extension types for OpenCL types on SPIR-V.
Joshua Cranmer via llvm-commits
llvm-commits at lists.llvm.org
Mon Mar 13 11:20:59 PDT 2023
Author: Joshua Cranmer
Date: 2023-03-13T14:20:24-04:00
New Revision: bcad161db3e69e27736c975ef5eeac60c96dcc97
URL: https://github.com/llvm/llvm-project/commit/bcad161db3e69e27736c975ef5eeac60c96dcc97
DIFF: https://github.com/llvm/llvm-project/commit/bcad161db3e69e27736c975ef5eeac60c96dcc97.diff
LOG: [Clang][SPIR-V] Emit target extension types for OpenCL types on SPIR-V.
Reviewed By: Anastasia
Differential Revision: https://reviews.llvm.org/D141008
Added:
Modified:
clang/include/clang-c/Index.h
clang/include/clang/Basic/OpenCLExtensionTypes.def
clang/lib/CodeGen/CGOpenCLRuntime.cpp
clang/lib/CodeGen/CGOpenCLRuntime.h
clang/lib/CodeGen/TargetInfo.cpp
clang/lib/CodeGen/TargetInfo.h
clang/test/CodeGenOpenCL/cast_image.cl
clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
clang/test/CodeGenOpenCL/opencl_types.cl
clang/test/CodeGenOpenCL/sampler.cl
clang/test/Index/pipe-size.cl
llvm/docs/SPIRVUsage.rst
Removed:
################################################################################
diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h
index 88d719d32c21f..5f71b86b20b48 100644
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@@ -2923,10 +2923,15 @@ enum CXTypeKind {
CXType_OCLIntelSubgroupAVCImeResult = 169,
CXType_OCLIntelSubgroupAVCRefResult = 170,
CXType_OCLIntelSubgroupAVCSicResult = 171,
+ CXType_OCLIntelSubgroupAVCImeResultSingleReferenceStreamout = 172,
+ CXType_OCLIntelSubgroupAVCImeResultDualReferenceStreamout = 173,
+ CXType_OCLIntelSubgroupAVCImeSingleReferenceStreamin = 174,
+ CXType_OCLIntelSubgroupAVCImeDualReferenceStreamin = 175,
+
+ /* Old aliases for AVC OpenCL extension types. */
CXType_OCLIntelSubgroupAVCImeResultSingleRefStreamout = 172,
CXType_OCLIntelSubgroupAVCImeResultDualRefStreamout = 173,
CXType_OCLIntelSubgroupAVCImeSingleRefStreamin = 174,
-
CXType_OCLIntelSubgroupAVCImeDualRefStreamin = 175,
CXType_ExtVector = 176,
diff --git a/clang/include/clang/Basic/OpenCLExtensionTypes.def b/clang/include/clang/Basic/OpenCLExtensionTypes.def
index 84ffbe936b77d..17c72d69a0206 100644
--- a/clang/include/clang/Basic/OpenCLExtensionTypes.def
+++ b/clang/include/clang/Basic/OpenCLExtensionTypes.def
@@ -28,10 +28,10 @@ INTEL_SUBGROUP_AVC_TYPE(mce_result_t, MceResult)
INTEL_SUBGROUP_AVC_TYPE(ime_result_t, ImeResult)
INTEL_SUBGROUP_AVC_TYPE(ref_result_t, RefResult)
INTEL_SUBGROUP_AVC_TYPE(sic_result_t, SicResult)
-INTEL_SUBGROUP_AVC_TYPE(ime_result_single_reference_streamout_t, ImeResultSingleRefStreamout)
-INTEL_SUBGROUP_AVC_TYPE(ime_result_dual_reference_streamout_t, ImeResultDualRefStreamout)
-INTEL_SUBGROUP_AVC_TYPE(ime_single_reference_streamin_t, ImeSingleRefStreamin)
-INTEL_SUBGROUP_AVC_TYPE(ime_dual_reference_streamin_t, ImeDualRefStreamin)
+INTEL_SUBGROUP_AVC_TYPE(ime_result_single_reference_streamout_t, ImeResultSingleReferenceStreamout)
+INTEL_SUBGROUP_AVC_TYPE(ime_result_dual_reference_streamout_t, ImeResultDualReferenceStreamout)
+INTEL_SUBGROUP_AVC_TYPE(ime_single_reference_streamin_t, ImeSingleReferenceStreamin)
+INTEL_SUBGROUP_AVC_TYPE(ime_dual_reference_streamin_t, ImeDualReferenceStreamin)
#undef INTEL_SUBGROUP_AVC_TYPE
#endif // INTEL_SUBGROUP_AVC_TYPE
diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.cpp b/clang/lib/CodeGen/CGOpenCLRuntime.cpp
index 5a5d7a4d18c57..395ba5e7dc144 100644
--- a/clang/lib/CodeGen/CGOpenCLRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenCLRuntime.cpp
@@ -31,8 +31,11 @@ void CGOpenCLRuntime::EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF,
}
llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
- assert(T->isOpenCLSpecificType() &&
- "Not an OpenCL specific type!");
+ assert(T->isOpenCLSpecificType() && "Not an OpenCL specific type!");
+
+ // Check if the target has a specific translation for this type first.
+ if (llvm::Type *TransTy = CGM.getTargetCodeGenInfo().getOpenCLType(CGM, T))
+ return TransTy;
switch (cast<BuiltinType>(T)->getKind()) {
default:
@@ -75,6 +78,9 @@ llvm::PointerType *CGOpenCLRuntime::getPointerType(const Type *T,
}
llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T) {
+ if (llvm::Type *PipeTy = CGM.getTargetCodeGenInfo().getOpenCLType(CGM, T))
+ return PipeTy;
+
if (T->isReadOnly())
return getPipeType(T, "opencl.pipe_ro_t", PipeROTy);
else
@@ -91,12 +97,18 @@ llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T, StringRef Name,
return PipeTy;
}
-llvm::PointerType *CGOpenCLRuntime::getSamplerType(const Type *T) {
- if (!SamplerTy)
- SamplerTy = llvm::PointerType::get(llvm::StructType::create(
- CGM.getLLVMContext(), "opencl.sampler_t"),
- CGM.getContext().getTargetAddressSpace(
- CGM.getContext().getOpenCLTypeAddrSpace(T)));
+llvm::Type *CGOpenCLRuntime::getSamplerType(const Type *T) {
+ if (SamplerTy)
+ return SamplerTy;
+
+ if (llvm::Type *TransTy = CGM.getTargetCodeGenInfo().getOpenCLType(
+ CGM, CGM.getContext().OCLSamplerTy.getTypePtr()))
+ SamplerTy = TransTy;
+ else
+ SamplerTy = llvm::PointerType::get(
+ llvm::StructType::create(CGM.getLLVMContext(), "opencl.sampler_t"),
+ CGM.getContext().getTargetAddressSpace(
+ CGM.getContext().getOpenCLTypeAddrSpace(T)));
return SamplerTy;
}
diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.h b/clang/lib/CodeGen/CGOpenCLRuntime.h
index b93e2041c8f16..df8084d6008be 100644
--- a/clang/lib/CodeGen/CGOpenCLRuntime.h
+++ b/clang/lib/CodeGen/CGOpenCLRuntime.h
@@ -38,7 +38,7 @@ class CGOpenCLRuntime {
CodeGenModule &CGM;
llvm::Type *PipeROTy;
llvm::Type *PipeWOTy;
- llvm::PointerType *SamplerTy;
+ llvm::Type *SamplerTy;
llvm::StringMap<llvm::PointerType *> CachedTys;
/// Structure for enqueued block information.
@@ -70,7 +70,7 @@ class CGOpenCLRuntime {
virtual llvm::Type *getPipeType(const PipeType *T);
- llvm::PointerType *getSamplerType(const Type *T);
+ llvm::Type *getSamplerType(const Type *T);
// Returns a value which indicates the size in bytes of the pipe
// element.
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index 307d1e78b99b1..ba7a51c8ba77c 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -10559,6 +10559,7 @@ class CommonSPIRTargetCodeGenInfo : public TargetCodeGenInfo {
}
unsigned getOpenCLKernelCallingConv() const override;
+ llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const override;
};
class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo {
public:
@@ -10973,6 +10974,80 @@ static bool getTypeString(SmallStringEnc &Enc, const Decl *D,
return false;
}
+/// Construct a SPIR-V target extension type for the given OpenCL image type.
+static llvm::Type *getSPIRVImageType(llvm::LLVMContext &Ctx, StringRef BaseType,
+ StringRef OpenCLName,
+ unsigned AccessQualifier) {
+ // These parameters compare to the operands of OpTypeImage (see
+ // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeImage
+ // for more details). The first 6 integer parameters all default to 0, and
+ // will be changed to 1 only for the image type(s) that set the parameter to
+ // one. The 7th integer parameter is the access qualifier, which is tacked on
+ // at the end.
+ SmallVector<unsigned, 7> IntParams = {0, 0, 0, 0, 0, 0};
+
+ // Choose the dimension of the image--this corresponds to the Dim enum in
+ // SPIR-V (first integer parameter of OpTypeImage).
+ if (OpenCLName.startswith("image2d"))
+ IntParams[0] = 1; // 1D
+ else if (OpenCLName.startswith("image3d"))
+ IntParams[0] = 2; // 2D
+ else if (OpenCLName == "image1d_buffer")
+ IntParams[0] = 5; // Buffer
+ else
+ assert(OpenCLName.startswith("image1d") && "Unknown image type");
+
+ // Set the other integer parameters of OpTypeImage if necessary. Note that the
+ // OpenCL image types don't provide any information for the Sampled or
+ // Image Format parameters.
+ if (OpenCLName.contains("_depth"))
+ IntParams[1] = 1;
+ if (OpenCLName.contains("_array"))
+ IntParams[2] = 1;
+ if (OpenCLName.contains("_msaa"))
+ IntParams[3] = 1;
+
+ // Access qualifier
+ IntParams.push_back(AccessQualifier);
+
+ return llvm::TargetExtType::get(Ctx, BaseType, {llvm::Type::getVoidTy(Ctx)},
+ IntParams);
+}
+
+llvm::Type *CommonSPIRTargetCodeGenInfo::getOpenCLType(CodeGenModule &CGM,
+ const Type *Ty) const {
+ llvm::LLVMContext &Ctx = CGM.getLLVMContext();
+ if (auto *PipeTy = dyn_cast<PipeType>(Ty))
+ return llvm::TargetExtType::get(Ctx, "spirv.Pipe", {},
+ {!PipeTy->isReadOnly()});
+ if (auto *BuiltinTy = dyn_cast<BuiltinType>(Ty)) {
+ enum AccessQualifier : unsigned { AQ_ro = 0, AQ_wo = 1, AQ_rw = 2 };
+ switch (BuiltinTy->getKind()) {
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
+ case BuiltinType::Id: \
+ return getSPIRVImageType(Ctx, "spirv.Image", #ImgType, AQ_##Suffix);
+#include "clang/Basic/OpenCLImageTypes.def"
+ case BuiltinType::OCLSampler:
+ return llvm::TargetExtType::get(Ctx, "spirv.Sampler");
+ case BuiltinType::OCLEvent:
+ return llvm::TargetExtType::get(Ctx, "spirv.Event");
+ case BuiltinType::OCLClkEvent:
+ return llvm::TargetExtType::get(Ctx, "spirv.DeviceEvent");
+ case BuiltinType::OCLQueue:
+ return llvm::TargetExtType::get(Ctx, "spirv.Queue");
+ case BuiltinType::OCLReserveID:
+ return llvm::TargetExtType::get(Ctx, "spirv.ReserveId");
+#define INTEL_SUBGROUP_AVC_TYPE(Name, Id) \
+ case BuiltinType::OCLIntelSubgroupAVC##Id: \
+ return llvm::TargetExtType::get(Ctx, "spirv.Avc" #Id "INTEL");
+#include "clang/Basic/OpenCLExtensionTypes.def"
+ default:
+ return nullptr;
+ }
+ }
+
+ return nullptr;
+}
//===----------------------------------------------------------------------===//
// RISCV ABI Implementation
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h
index b97c068bb9b35..b7e6ae85be203 100644
--- a/clang/lib/CodeGen/TargetInfo.h
+++ b/clang/lib/CodeGen/TargetInfo.h
@@ -379,6 +379,11 @@ class TargetCodeGenInfo {
// DO NOTHING by default.
return false;
}
+
+ /// Return an LLVM type that corresponds to an OpenCL type.
+ virtual llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const {
+ return nullptr;
+ }
};
} // namespace CodeGen
diff --git a/clang/test/CodeGenOpenCL/cast_image.cl b/clang/test/CodeGenOpenCL/cast_image.cl
index af68a7e720742..51f46fa7712ed 100644
--- a/clang/test/CodeGenOpenCL/cast_image.cl
+++ b/clang/test/CodeGenOpenCL/cast_image.cl
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s
-// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s
+// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple x86_64-unknown-unknown %s | FileCheck --check-prefix=X86 %s
#ifdef __AMDGCN__
@@ -11,7 +11,7 @@ constant int* convert(image2d_t img) {
#else
global int* convert(image2d_t img) {
- // SPIR: bitcast %opencl.image2d_ro_t addrspace(1)* %img to i32 addrspace(1)*
+ // X86: bitcast %opencl.image2d_ro_t* %img to i32*
return __builtin_astype(img, global int*);
}
diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
index a917d128d09d0..88fafbbc98f43 100644
--- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
+++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
@@ -27,7 +27,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// SPIR32-NEXT: [[A_ADDR:%.*]] = alloca ptr addrspace(1), align 4
// SPIR32-NEXT: [[B_ADDR:%.*]] = alloca ptr addrspace(1), align 4
// SPIR32-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4
-// SPIR32-NEXT: [[DEFAULT_QUEUE:%.*]] = alloca ptr, align 4
+// SPIR32-NEXT: [[DEFAULT_QUEUE:%.*]] = alloca target("spirv.Queue"), align 4
// SPIR32-NEXT: [[FLAGS:%.*]] = alloca i32, align 4
// SPIR32-NEXT: [[NDRANGE:%.*]] = alloca [[STRUCT_NDRANGE_T:%.*]], align 4
// SPIR32-NEXT: [[TMP:%.*]] = alloca [[STRUCT_NDRANGE_T]], align 4
@@ -36,7 +36,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// SPIR32-NEXT: store ptr addrspace(1) [[B]], ptr [[B_ADDR]], align 4
// SPIR32-NEXT: store i32 [[I]], ptr [[I_ADDR]], align 4
// SPIR32-NEXT: store i32 0, ptr [[FLAGS]], align 4
-// SPIR32-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DEFAULT_QUEUE]], align 4
+// SPIR32-NEXT: [[TMP0:%.*]] = load target("spirv.Queue"), ptr [[DEFAULT_QUEUE]], align 4
// SPIR32-NEXT: [[TMP1:%.*]] = load i32, ptr [[FLAGS]], align 4
// SPIR32-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP]], ptr align 4 [[NDRANGE]], i32 4, i1 false)
// SPIR32-NEXT: [[BLOCK_SIZE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 0
@@ -55,7 +55,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// SPIR32-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[B_ADDR]], align 4
// SPIR32-NEXT: store ptr addrspace(1) [[TMP4]], ptr [[BLOCK_CAPTURED2]], align 4
// SPIR32-NEXT: [[TMP5:%.*]] = addrspacecast ptr [[BLOCK]] to ptr addrspace(4)
-// SPIR32-NEXT: [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(ptr [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
+// SPIR32-NEXT: [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
// SPIR32-NEXT: ret void
//
//
@@ -98,7 +98,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// STRICTFP-NEXT: [[A_ADDR:%.*]] = alloca ptr addrspace(1), align 4
// STRICTFP-NEXT: [[B_ADDR:%.*]] = alloca ptr addrspace(1), align 4
// STRICTFP-NEXT: [[I_ADDR:%.*]] = alloca i32, align 4
-// STRICTFP-NEXT: [[DEFAULT_QUEUE:%.*]] = alloca ptr, align 4
+// STRICTFP-NEXT: [[DEFAULT_QUEUE:%.*]] = alloca target("spirv.Queue"), align 4
// STRICTFP-NEXT: [[FLAGS:%.*]] = alloca i32, align 4
// STRICTFP-NEXT: [[NDRANGE:%.*]] = alloca [[STRUCT_NDRANGE_T:%.*]], align 4
// STRICTFP-NEXT: [[TMP:%.*]] = alloca [[STRUCT_NDRANGE_T]], align 4
@@ -107,7 +107,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// STRICTFP-NEXT: store ptr addrspace(1) [[B]], ptr [[B_ADDR]], align 4
// STRICTFP-NEXT: store i32 [[I]], ptr [[I_ADDR]], align 4
// STRICTFP-NEXT: store i32 0, ptr [[FLAGS]], align 4
-// STRICTFP-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DEFAULT_QUEUE]], align 4
+// STRICTFP-NEXT: [[TMP0:%.*]] = load target("spirv.Queue"), ptr [[DEFAULT_QUEUE]], align 4
// STRICTFP-NEXT: [[TMP1:%.*]] = load i32, ptr [[FLAGS]], align 4
// STRICTFP-NEXT: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP]], ptr align 4 [[NDRANGE]], i32 4, i1 false) #[[ATTR5:[0-9]+]]
// STRICTFP-NEXT: [[BLOCK_SIZE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 0
@@ -126,7 +126,7 @@ kernel void device_side_enqueue(global float *a, global float *b, int i) {
// STRICTFP-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[B_ADDR]], align 4
// STRICTFP-NEXT: store ptr addrspace(1) [[TMP4]], ptr [[BLOCK_CAPTURED2]], align 4
// STRICTFP-NEXT: [[TMP5:%.*]] = addrspacecast ptr [[BLOCK]] to ptr addrspace(4)
-// STRICTFP-NEXT: [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(ptr [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
+// STRICTFP-NEXT: [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
// STRICTFP-NEXT: ret void
//
//
diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
index 05a15e476f211..f0c164795b764 100644
--- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
+++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
@@ -1,9 +1,12 @@
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR
// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR
// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefix=CHECK-LIFETIMES
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
@@ -29,7 +32,7 @@ typedef struct {int a;} ndrange_t;
// COMMON: [[BLG11:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr {{@[^ ]+}} to ptr addrspace(4)) }
// Emits block literal [[BL_GLOBAL]], invoke function [[INV_G]] and global block variable @block_G
-// COMMON: define internal spir_func void [[INV_G]](ptr addrspace(4) %{{.*}}, ptr addrspace(3) %{{.*}})
+// COMMON: define internal {{(spir_func )?}}void [[INV_G]](ptr addrspace(4) %{{.*}}, ptr addrspace(3) %{{.*}})
const bl_t block_G = (bl_t) ^ (local void *a) {};
void callee(int id, __global int *out) {
@@ -38,17 +41,21 @@ void callee(int id, __global int *out) {
// COMMON-LABEL: define{{.*}} spir_kernel void @device_side_enqueue(ptr addrspace(1) align 4 %{{.*}}, ptr addrspace(1) align 4 %b, i32 %i)
kernel void device_side_enqueue(global int *a, global int *b, int i) {
- // COMMON: %default_queue = alloca ptr
+ // SPIR: %default_queue = alloca target("spirv.Queue")
+ // X86: %default_queue = alloca ptr
queue_t default_queue;
// COMMON: %flags = alloca i32
unsigned flags = 0;
// COMMON: %ndrange = alloca %struct.ndrange_t
ndrange_t ndrange;
- // COMMON: %clk_event = alloca ptr
+ // SPIR: %clk_event = alloca target("spirv.DeviceEvent")
+ // X86: %clk_event = alloca ptr
clk_event_t clk_event;
- // COMMON: %event_wait_list = alloca ptr
+ // SPIR: %event_wait_list = alloca target("spirv.DeviceEvent")
+ // X86: %event_wait_list = alloca ptr
clk_event_t event_wait_list;
- // COMMON: %event_wait_list2 = alloca [1 x ptr]
+ // SPIR: %event_wait_list2 = alloca [1 x target("spirv.DeviceEvent")]
+ // X86: %event_wait_list2 = alloca [1 x ptr]
clk_event_t event_wait_list2[] = {clk_event};
// COMMON: [[NDR:%[a-z0-9]+]] = alloca %struct.ndrange_t, align 4
@@ -76,12 +83,14 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// CHECK-LIFETIMES: %[[BLOCK_SIZES7:.*]] = alloca [1 x i64]
// Emits block literal on stack and block kernel [[INVLK1]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL1:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
// COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr %block to ptr addrspace(4)
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
- // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic(
+ // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK1:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange,
@@ -90,14 +99,16 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
});
// Emits block literal on stack and block kernel [[INVLK2]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %event_wait_list to ptr addrspace(4)
// COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
// COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL2:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
// COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr %block4 to ptr addrspace(4)
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
- // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic_events
+ // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+ // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK2:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
@@ -105,25 +116,28 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
a[i] = b[i];
});
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
- // COMMON-SAME: (ptr {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic_events
+ // SPIR-SAME: (target("spirv.Queue") {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
+ // X86-SAME: (ptr {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
enqueue_kernel(default_queue, flags, ndrange, 1, 0, 0,
^(void) {
return;
});
// Emits global block literal [[BLG1]] and block kernel [[INVGK1]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// CHECK-LIFETIMES: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES1]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES1]])
// B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES1]], i32 0, i32 0
// B32: store i32 256, ptr %[[TMP]], align 4
// B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES1]], i32 0, i32 0
// B64: store i64 256, ptr %[[TMP]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
- // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
+ // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK1:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG1]] to ptr addrspace(4)), i32 1,
// B32-SAME: ptr %[[TMP]])
@@ -136,17 +150,19 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
char c;
// Emits global block literal [[BLG2]] and block kernel [[INVGK2]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES2]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES2]])
// B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES2]], i32 0, i32 0
// B32: store i32 %{{.*}}, ptr %[[TMP]], align 4
// B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES2]], i32 0, i32 0
// B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
- // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
+ // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK2:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG2]] to ptr addrspace(4)), i32 1,
// B32-SAME: ptr %[[TMP]])
@@ -158,20 +174,23 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
c);
// Emits global block literal [[BLG3]] and block kernel [[INVGK3]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
- // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+ // SPIR: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+ // X86: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
// COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr [[AD]] to ptr addrspace(4)
// COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
// CHECK-LIFETIMES: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES3]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES3]])
// B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES3]], i32 0, i32 0
// B32: store i32 256, ptr %[[TMP]], align 4
// B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES3]], i32 0, i32 0
// B64: store i64 256, ptr %[[TMP]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
- // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs
+ // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+ // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK3:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG3]] to ptr addrspace(4)), i32 1,
// B32-SAME: ptr %[[TMP]])
@@ -183,20 +202,23 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
256);
// Emits global block literal [[BLG4]] and block kernel [[INVGK4]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
- // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+ // SPIR: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+ // X86: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
// COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr [[AD]] to ptr addrspace(4)
// COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES4]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES4]])
// B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES4]], i32 0, i32 0
// B32: store i32 %{{.*}}, ptr %[[TMP]], align 4
// B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES4]], i32 0, i32 0
// B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
- // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs
+ // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+ // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK4:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG4]] to ptr addrspace(4)), i32 1,
// B32-SAME: ptr %[[TMP]])
@@ -209,17 +231,19 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
long l;
// Emits global block literal [[BLG5]] and block kernel [[INVGK5]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES5]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES5]])
// B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES5]], i32 0, i32 0
// B32: store i32 %{{.*}}, ptr %[[TMP]], align 4
// B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES5]], i32 0, i32 0
// B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
- // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs
+ // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK5:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG5]] to ptr addrspace(4)), i32 1,
// B32-SAME: ptr %[[TMP]])
@@ -231,10 +255,11 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
l);
// Emits global block literal [[BLG6]] and block kernel [[INVGK6]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %[[BLOCK_SIZES6]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %[[BLOCK_SIZES6]])
// B32: %[[TMP:.*]] = getelementptr [3 x i32], ptr %[[BLOCK_SIZES6]], i32 0, i32 0
// B32: store i32 1, ptr %[[TMP]], align 4
@@ -248,8 +273,9 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// B64: store i64 2, ptr %[[BLOCK_SIZES62]], align 8
// B64: %[[BLOCK_SIZES63:.*]] = getelementptr [3 x i64], ptr %[[BLOCK_SIZES6]], i32 0, i32 2
// B64: store i64 4, ptr %[[BLOCK_SIZES63]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
- // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs
+ // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK6:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG6]] to ptr addrspace(4)), i32 3,
// B32-SAME: ptr %[[TMP]])
@@ -261,17 +287,19 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
1, 2, 4);
// Emits global block literal [[BLG7]] and block kernel [[INVGK7]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES7]])
- // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+ // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES7]])
// B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES7]], i32 0, i32 0
// B32: store i32 0, ptr %[[TMP]], align 4
// B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES7]], i32 0, i32 0
// B64: store i64 4294967296, ptr %[[TMP]], align 8
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
- // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs
+ // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK7:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG7]] to ptr addrspace(4)), i32 1,
// B32-SAME: ptr %[[TMP]])
@@ -297,26 +325,28 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
};
// Uses global block literal [[BLG8]] and invoke function [[INVG8]].
- // COMMON: call spir_func void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4))) [[INVOKE_ATTR:#[0-9]+]]
+ // COMMON: call {{(spir_func )?}}void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4))) [[INVOKE_ATTR:#[0-9]+]]
block_A();
// Emits global block literal [[BLG8]] and block kernel [[INVGK8]]. [[INVGK8]] calls [[INVG8]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
- // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic(
+ // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
enqueue_kernel(default_queue, flags, ndrange, block_A);
// Uses block kernel [[INVGK8]] and global block literal [[BLG8]].
- // COMMON: call spir_func i32 @__get_kernel_work_group_size_impl(
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_work_group_size_impl(
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
unsigned size = get_kernel_work_group_size(block_A);
// Uses global block literal [[BLG8]] and invoke function [[INVG8]]. Make sure no redundant block literal and invoke functions are emitted.
- // COMMON: call spir_func void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
+ // COMMON: call {{(spir_func )?}}void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
block_A();
// Make sure that block invoke function is resolved correctly after sequence of assignements.
@@ -332,12 +362,12 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON-SAME: to ptr addrspace(4)),
// COMMON-SAME: ptr %b2,
bl_t b2 = b1;
- // COMMON: call spir_func void @block_G_block_invoke(ptr addrspace(4) addrspacecast (ptr addrspace(1)
+ // COMMON: call {{(spir_func )?}}void @block_G_block_invoke(ptr addrspace(4) addrspacecast (ptr addrspace(1)
// COMMON-SAME: [[BL_GLOBAL]]
// COOMON-SAME: to ptr addrspace(4)), ptr addrspace(3) null)
b2(0);
// Uses global block literal [[BL_GLOBAL]] and block kernel [[INV_G_K]]. [[INV_G_K]] calls [[INV_G]].
- // COMMON: call spir_func i32 @__get_kernel_preferred_work_group_size_multiple_impl(
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_preferred_work_group_size_multiple_impl(
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INV_G_K:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BL_GLOBAL]] to ptr addrspace(4)))
size = get_kernel_preferred_work_group_size_multiple(b2);
@@ -347,41 +377,43 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
};
// Emits block literal on stack and block kernel [[INVLK3]].
// COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL3:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
- // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+ // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+ // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
// COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr {{.*}} to ptr addrspace(4)
- // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
- // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic(
+ // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK3:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange, block_C);
// Emits global block literal [[BLG9]] and block kernel [[INVGK9]]. [[INVGK9]] calls [[INV9]].
- // COMMON: call spir_func i32 @__get_kernel_work_group_size_impl(
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_work_group_size_impl(
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK9:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG9]] to ptr addrspace(4)))
size = get_kernel_work_group_size(block_B);
// Uses global block literal [[BLG8]] and block kernel [[INVGK8]]. Make sure no redundant block literal ind invoke functions are emitted.
- // COMMON: call spir_func i32 @__get_kernel_preferred_work_group_size_multiple_impl(
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_preferred_work_group_size_multiple_impl(
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
size = get_kernel_preferred_work_group_size_multiple(block_A);
// Uses global block literal [[BL_GLOBAL]] and block kernel [[INV_G_K]]. [[INV_G_K]] calls [[INV_G]].
- // COMMON: call spir_func i32 @__get_kernel_preferred_work_group_size_multiple_impl(
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_preferred_work_group_size_multiple_impl(
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INV_G_K:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BL_GLOBAL]] to ptr addrspace(4)))
size = get_kernel_preferred_work_group_size_multiple(block_G);
// Emits global block literal [[BLG10]] and block kernel [[INVGK10]].
- // COMMON: call spir_func i32 @__get_kernel_max_sub_group_size_for_ndrange_impl(ptr {{[^,]+}},
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_max_sub_group_size_for_ndrange_impl(ptr {{[^,]+}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK10:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG10]] to ptr addrspace(4)))
size = get_kernel_max_sub_group_size_for_ndrange(ndrange, ^(){});
// Emits global block literal [[BLG11]] and block kernel [[INVGK11]].
- // COMMON: call spir_func i32 @__get_kernel_sub_group_count_for_ndrange_impl(ptr {{[^,]+}},
+ // COMMON: call {{(spir_func )?}}i32 @__get_kernel_sub_group_count_for_ndrange_impl(ptr {{[^,]+}},
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK11:[^ ]+_kernel]] to ptr addrspace(4)),
// COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG11]] to ptr addrspace(4)))
size = get_kernel_sub_group_count_for_ndrange(ndrange, ^(){});
@@ -389,7 +421,7 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: define spir_kernel void [[INVLK1]](ptr addrspace(4) %0) #{{[0-9]+}} {
// COMMON: entry:
-// COMMON: call spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) %0)
+// COMMON: call {{(spir_func )?}}void @__device_side_enqueue_block_invoke(ptr addrspace(4) %0)
// COMMON: ret void
// COMMON: }
// COMMON: define spir_kernel void [[INVLK2]](ptr addrspace(4){{.*}})
@@ -400,12 +432,12 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: define spir_kernel void [[INVGK5]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
// COMMON: define spir_kernel void [[INVGK6]](ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3) #{{[0-9]+}} {
// COMMON: entry:
-// COMMON: call spir_func void @__device_side_enqueue_block_invoke_9(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3)
+// COMMON: call {{(spir_func )?}}void @__device_side_enqueue_block_invoke_9(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3)
// COMMON: ret void
// COMMON: }
// COMMON: define spir_kernel void [[INVGK7]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
-// COMMON: define internal spir_func void [[INVG8]](ptr addrspace(4){{.*}}) [[INVG8_INVOKE_FUNC_ATTR:#[0-9]+]]
-// COMMON: define internal spir_func void [[INVG9]](ptr addrspace(4){{.*}}, ptr addrspace(3) %{{.*}})
+// COMMON: define internal {{(spir_func )?}}void [[INVG8]](ptr addrspace(4){{.*}}) [[INVG8_INVOKE_FUNC_ATTR:#[0-9]+]]
+// COMMON: define internal {{(spir_func )?}}void [[INVG9]](ptr addrspace(4){{.*}}, ptr addrspace(3) %{{.*}})
// COMMON: define spir_kernel void [[INVGK8]](ptr addrspace(4){{.*}})
// COMMON: define spir_kernel void [[INV_G_K]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
// COMMON: define spir_kernel void [[INVLK3]](ptr addrspace(4){{.*}})
@@ -413,5 +445,7 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: define spir_kernel void [[INVGK10]](ptr addrspace(4){{.*}})
// COMMON: define spir_kernel void [[INVGK11]](ptr addrspace(4){{.*}})
-// COMMON: attributes [[INVG8_INVOKE_FUNC_ATTR]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// COMMON: attributes [[INVOKE_KERNEL_ATTR]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// SPIR: attributes [[INVG8_INVOKE_FUNC_ATTR]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// SPIR: attributes [[INVOKE_KERNEL_ATTR]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// X86: attributes [[INVG8_INVOKE_FUNC_ATTR]] = { convergent noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="{{[^"]*}}" }
+// X86: attributes [[INVOKE_KERNEL_ATTR]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="{{[^"]*}}" }
diff --git a/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl b/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
index 3c55276213e62..49e1eaf91fdd2 100644
--- a/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
+++ b/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
@@ -1,45 +1,30 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -emit-llvm -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -emit-llvm -o - -O0 | FileCheck %s
-// CHECK: %opencl.intel_sub_group_avc_mce_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ref_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_sic_payload_t = type opaque
+// CHECK: store target("spirv.AvcImePayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefPayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicPayloadINTEL") zeroinitializer,
-// CHECK: %opencl.intel_sub_group_avc_mce_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ref_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_sic_result_t = type opaque
+// CHECK: store target("spirv.AvcImeResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicResultINTEL") zeroinitializer,
-// CHECK: %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_single_reference_streamin_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t = type opaque
+// CHECK: store target("spirv.AvcImeResultSingleReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeResultDualReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeSingleReferenceStreaminINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeDualReferenceStreaminINTEL") zeroinitializer,
-// CHECK: store %opencl.intel_sub_group_avc_ime_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_payload_t* null,
+// CHECK: store target("spirv.AvcImePayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefPayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicPayloadINTEL") zeroinitializer,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_result_t* null,
+// CHECK: store target("spirv.AvcImeResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicResultINTEL") zeroinitializer,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_single_reference_streamin_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t* null,
-//
-// CHECK: store %opencl.intel_sub_group_avc_ime_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_payload_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_result_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_single_reference_streamin_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t* null,
+// CHECK: store target("spirv.AvcImeResultSingleReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeResultDualReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeSingleReferenceStreaminINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeDualReferenceStreaminINTEL") zeroinitializer,
#pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable
diff --git a/clang/test/CodeGenOpenCL/opencl_types.cl b/clang/test/CodeGenOpenCL/opencl_types.cl
index 17e70a9f774fb..b8ccfa26f16d9 100644
--- a/clang/test/CodeGenOpenCL/opencl_types.cl
+++ b/clang/test/CodeGenOpenCL/opencl_types.cl
@@ -10,65 +10,65 @@ constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRU
// CHECK-COM-NOT: constant i32
void fnc1(image1d_t img) {}
-// CHECK-SPIR: @fnc1(ptr addrspace(1)
+// CHECK-SPIR: @fnc1(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0)
// CHECK-AMDGCN: @fnc1(ptr addrspace(4)
void fnc1arr(image1d_array_t img) {}
-// CHECK-SPIR: @fnc1arr(ptr addrspace(1)
+// CHECK-SPIR: @fnc1arr(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0)
// CHECK-AMDGCN: @fnc1arr(ptr addrspace(4)
void fnc1buff(image1d_buffer_t img) {}
-// CHECK-SPIR: @fnc1buff(ptr addrspace(1)
+// CHECK-SPIR: @fnc1buff(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0)
// CHECK-AMDGCN: @fnc1buff(ptr addrspace(4)
void fnc2(image2d_t img) {}
-// CHECK-SPIR: @fnc2(ptr addrspace(1)
+// CHECK-SPIR: @fnc2(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)
// CHECK-AMDGCN: @fnc2(ptr addrspace(4)
void fnc2arr(image2d_array_t img) {}
-// CHECK-SPIR: @fnc2arr(ptr addrspace(1)
+// CHECK-SPIR: @fnc2arr(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0)
// CHECK-AMDGCN: @fnc2arr(ptr addrspace(4)
void fnc3(image3d_t img) {}
-// CHECK-SPIR: @fnc3(ptr addrspace(1)
+// CHECK-SPIR: @fnc3(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0)
// CHECK-AMDGCN: @fnc3(ptr addrspace(4)
void fnc4smp(sampler_t s) {}
-// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(2)
+// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(target("spirv.Sampler")
// CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(4)
kernel void foo(image1d_t img) {
sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR;
- // CHECK-SPIR: alloca ptr addrspace(2)
+ // CHECK-SPIR: alloca target("spirv.Sampler")
// CHECK-AMDGCN: alloca ptr addrspace(4)
event_t evt;
- // CHECK-SPIR: alloca ptr
+ // CHECK-SPIR: alloca target("spirv.Event")
// CHECK-AMDGCN: alloca ptr addrspace(5)
clk_event_t clk_evt;
- // CHECK-SPIR: alloca ptr
+ // CHECK-SPIR: alloca target("spirv.DeviceEvent")
// CHECK-AMDGCN: alloca ptr addrspace(1)
queue_t queue;
- // CHECK-SPIR: alloca ptr
+ // CHECK-SPIR: alloca target("spirv.Queue")
// CHECK-AMDGCN: alloca ptr addrspace(1)
reserve_id_t rid;
- // CHECK-SPIR: alloca ptr
+ // CHECK-SPIR: alloca target("spirv.ReserveId")
// CHECK-AMDGCN: alloca ptr addrspace(1)
- // CHECK-SPIR: store ptr addrspace(2)
+ // CHECK-SPIR: store target("spirv.Sampler")
// CHECK-AMDGCN: store ptr addrspace(4)
fnc4smp(smp);
- // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2)
+ // CHECK-SPIR: call {{.*}}void @fnc4smp(target("spirv.Sampler")
// CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4)
fnc4smp(glb_smp);
- // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2)
+ // CHECK-SPIR: call {{.*}}void @fnc4smp(target("spirv.Sampler")
// CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4)
}
kernel void foo_ro_pipe(read_only pipe int p) {}
-// CHECK-SPIR: @foo_ro_pipe(ptr addrspace(1) %p)
+// CHECK-SPIR: @foo_ro_pipe(target("spirv.Pipe", 0) %p)
// CHECK_AMDGCN: @foo_ro_pipe(ptr addrspace(1) %p)
kernel void foo_wo_pipe(write_only pipe int p) {}
-// CHECK-SPIR: @foo_wo_pipe(ptr addrspace(1) %p)
+// CHECK-SPIR: @foo_wo_pipe(target("spirv.Pipe", 1) %p)
// CHECK_AMDGCN: @foo_wo_pipe(ptr addrspace(1) %p)
void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
diff --git a/clang/test/CodeGenOpenCL/sampler.cl b/clang/test/CodeGenOpenCL/sampler.cl
index 6c367a2327e37..3649afba50fe2 100644
--- a/clang/test/CodeGenOpenCL/sampler.cl
+++ b/clang/test/CodeGenOpenCL/sampler.cl
@@ -1,6 +1,9 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-SPIR %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-SPIR %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-SPIR %s
+// RUN: %clang_cc1 %s -emit-llvm -triple x86_64-unknown-linux-gnu -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-X86 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-X86 %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple x86_64-unknown-linux-gnu -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-X86 %s
//
// This test covers 5 cases of sampler initialzation:
// 1. function argument passing
@@ -17,74 +20,98 @@
#define CLK_FILTER_NEAREST 0x10
#define CLK_FILTER_LINEAR 0x20
-// CHECK: %opencl.sampler_t = type opaque
-
// Case 2a
constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
-// CHECK-NOT: glb_smp
+// CHECK-COMMON-NOT: glb_smp
// Case 2c
const sampler_t glb_smp_const = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
-// CHECK-NOT: glb_smp_const
+// CHECK-COMMON-NOT: glb_smp_const
int get_sampler_initializer(void);
void fnc4smp(sampler_t s) {}
-// CHECK: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](%opencl.sampler_t addrspace(2)* %
+// CHECK-SPIR: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](target("spirv.Sampler") %
+// CHECK-X86: define{{.*}} void [[FUNCNAME:@.*fnc4smp.*]](ptr %
kernel void foo(sampler_t smp_par) {
- // CHECK-LABEL: define{{.*}} spir_kernel void @foo(%opencl.sampler_t addrspace(2)* %smp_par)
- // CHECK: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)*
+ // CHECK-SPIR-LABEL: define{{.*}} spir_kernel void @foo(target("spirv.Sampler") %smp_par)
+ // CHECK-SPIR: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
+ // CHECK-X86-LABEL: define{{.*}} spir_kernel void @foo(ptr %smp_par)
+ // CHECK-X86: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca ptr
// Case 2b
sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST;
- // CHECK: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)*
- // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
- // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMP]], %opencl.sampler_t addrspace(2)** [[smp_ptr]]
+ // CHECK-SPIR: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+ // CHECK-SPIR: store target("spirv.Sampler") [[SAMP]], ptr [[smp_ptr]]
+ // CHECK-X86: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca ptr
+ // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 19)
+ // CHECK-X86: store ptr [[SAMP]], ptr [[smp_ptr]]
// Case 1b
fnc4smp(smp);
- // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
- // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]]
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86-NOT: call ptr @__translate_sampler_initializer(i32 19)
+ // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[smp_ptr]]
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
// Case 1b
fnc4smp(smp);
- // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
- // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]]
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86-NOT: call ptr @__translate_sampler_initializer(i32 19)
+ // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[smp_ptr]]
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
// Case 1a/2a
fnc4smp(glb_smp);
- // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
// Case 1a/2c
fnc4smp(glb_smp_const);
- // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
// Case 1c
fnc4smp(smp_par);
- // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_par_ptr]]
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_par_ptr]]
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[smp_par_ptr]]
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
fnc4smp(5);
- // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 5)
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 5)
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 5)
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
const sampler_t const_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
fnc4smp(const_smp);
- // CHECK: [[CONST_SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
- // CHECK: store %opencl.sampler_t addrspace(2)* [[CONST_SAMP]], %opencl.sampler_t addrspace(2)** [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
+ // CHECK-SPIR: [[CONST_SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK-SPIR: store target("spirv.Sampler") [[CONST_SAMP]], ptr [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
+ // CHECK-X86: [[CONST_SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+ // CHECK-X86: store ptr [[CONST_SAMP]], ptr [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
fnc4smp(const_smp);
- // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[CONST_SMP_PTR]]
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[CONST_SMP_PTR]]
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[CONST_SMP_PTR]]
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
constant sampler_t constant_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
fnc4smp(constant_smp);
- // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
- // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+ // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+ // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+ // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
// TODO: enable sampler initialization with non-constant integer.
//const sampler_t const_smp_func_init = get_sampler_initializer();
diff --git a/clang/test/Index/pipe-size.cl b/clang/test/Index/pipe-size.cl
index 0a384db00ba66..08b936f1a9b07 100644
--- a/clang/test/Index/pipe-size.cl
+++ b/clang/test/Index/pipe-size.cl
@@ -7,9 +7,9 @@ __kernel void testPipe( pipe int test )
int s = sizeof(test);
// X86: store ptr %test, ptr %test.addr, align 8
// X86: store i32 8, ptr %s, align 4
- // SPIR: store ptr addrspace(1) %test, ptr %test.addr, align 4
+ // SPIR: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 4
// SPIR: store i32 4, ptr %s, align 4
- // SPIR64: store ptr addrspace(1) %test, ptr %test.addr, align 8
+ // SPIR64: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 8
// SPIR64: store i32 8, ptr %s, align 4
// AMDGCN: store ptr addrspace(1) %test, ptr addrspace(5) %test.addr, align 8
// AMDGCN: store i32 8, ptr addrspace(5) %s, align 4
diff --git a/llvm/docs/SPIRVUsage.rst b/llvm/docs/SPIRVUsage.rst
index 0f75e001e661e..5ca035b68392f 100644
--- a/llvm/docs/SPIRVUsage.rst
+++ b/llvm/docs/SPIRVUsage.rst
@@ -75,3 +75,36 @@ to specify the target triple:
Example:
``-target spirv64v1.0`` can be used to compile for SPIR-V version 1.0 with 64-bit pointer width.
+
+.. _spirv-types:
+
+Representing special types in SPIR-V
+====================================
+
+SPIR-V specifies several kinds of opaque types. These types are represented
+using target extension types. These types are represented as follows:
+
+ .. table:: SPIR-V Opaque Types
+
+ ================== ====================== =========================================================================================
+ SPIR-V Type LLVM type name LLVM type arguments
+ ================== ====================== =========================================================================================
+ OpTypeImage ``spirv.Image`` sampled type, dimensionality, depth, arrayed, MS, sampled, image format, access qualifier
+ OpTypeSampler ``spirv.Sampler`` (none)
+ OpTypeSampledImage ``spirv.SampledImage`` sampled type, dimensionality, depth, arrayed, MS, sampled, image format, access qualifier
+ OpTypeEvent ``spirv.Event`` (none)
+ OpTypeDeviceEvent ``spirv.DeviceEvent`` (none)
+ OpTypeReserveId ``spirv.ReserveId`` (none)
+ OpTypeQueue ``spirv.Queue`` (none)
+ OpTypePipe ``spirv.Pipe`` access qualifier
+ OpTypePipeStorage ``spirv.PipeStorage`` (none)
+ ================== ====================== =========================================================================================
+
+All integer arguments take the same value as they do in their `corresponding
+SPIR-V instruction <https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_type_declaration_instructions>`_.
+For example, the OpenCL type ``image2d_depth_ro_t`` would be represented in
+SPIR-V IR as ``target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)``, with its
+dimensionality parameter as ``1`` meaning 2D. Sampled image types include the
+parameters of its underlying image type, so that a sampled image for the
+previous type has the representation
+``target("spirv.SampledImage, void, 1, 1, 0, 0, 0, 0, 0)``.
More information about the llvm-commits
mailing list