jcranmer-intel updated this revision to Diff 487960.
jcranmer-intel added a comment.
Fix some of the code review comments.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D141008/new/
https://reviews.llvm.org/D141008
Files:
clang/include/clang-c/Index.h
clang/include/clang/Basic/OpenCLExtensionTypes.def
clang/lib/CodeGen/CGExprScalar.cpp
clang/lib/CodeGen/CGOpenCLRuntime.cpp
clang/lib/CodeGen/CGOpenCLRuntime.h
clang/lib/CodeGen/CodeGenModule.cpp
clang/test/CodeGenOpenCL/cast_image.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
Index: llvm/docs/SPIRVUsage.rst
===================================================================
--- llvm/docs/SPIRVUsage.rst
+++ llvm/docs/SPIRVUsage.rst
@@ -75,3 +75,35 @@
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 the SPIR-V instruction.
+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)``.
Index: clang/test/Index/pipe-size.cl
===================================================================
--- clang/test/Index/pipe-size.cl
+++ clang/test/Index/pipe-size.cl
@@ -1,16 +1,16 @@
// RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
-// RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
-// RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
+// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
+// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
// RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa %s -o - | FileCheck %s --check-prefix=AMDGCN
__kernel void testPipe( pipe int test )
{
int s = sizeof(test);
// X86: store %opencl.pipe_ro_t* %test, %opencl.pipe_ro_t** %test.addr, align 8
// X86: store i32 8, i32* %s, align 4
- // SPIR: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)** %test.addr, align 4
- // SPIR: store i32 4, i32* %s, align 4
- // SPIR64: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)** %test.addr, align 8
- // SPIR64: store i32 8, i32* %s, align 4
+ // SPIR: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 4
+ // SPIR: store i32 4, ptr %s, align 4
+ // SPIR64: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 8
+ // SPIR64: store i32 8, ptr %s, align 4
// AMDGCN: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)* addrspace(5)* %test.addr, align 8
// AMDGCN: store i32 8, i32 addrspace(5)* %s, align 4
}
Index: clang/test/CodeGenOpenCL/sampler.cl
===================================================================
--- clang/test/CodeGenOpenCL/sampler.cl
+++ clang/test/CodeGenOpenCL/sampler.cl
@@ -1,6 +1,6 @@
-// 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 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
//
// This test covers 5 cases of sampler initialzation:
// 1. function argument passing
@@ -17,8 +17,6 @@
#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
@@ -30,61 +28,61 @@
int get_sampler_initializer(void);
void fnc4smp(sampler_t s) {}
-// CHECK: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](%opencl.sampler_t addrspace(2)* %
+// CHECK: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](target("spirv.Sampler") %
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-LABEL: define{{.*}} spir_kernel void @foo(target("spirv.Sampler") %smp_par)
+ // CHECK: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
// 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: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
+ // CHECK: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+ // CHECK: store target("spirv.Sampler") [[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-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+ // CHECK: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+ // CHECK: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_par_ptr]]
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 5)
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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: [[CONST_SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK: store target("spirv.Sampler") [[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: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[CONST_SMP_PTR]]
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[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: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+ // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
// TODO: enable sampler initialization with non-constant integer.
//const sampler_t const_smp_func_init = get_sampler_initializer();
Index: clang/test/CodeGenOpenCL/opencl_types.cl
===================================================================
--- clang/test/CodeGenOpenCL/opencl_types.cl
+++ clang/test/CodeGenOpenCL/opencl_types.cl
@@ -10,65 +10,65 @@
// 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) {}
Index: clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
===================================================================
--- clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
+++ 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
Index: clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
===================================================================
--- clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
+++ clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
@@ -39,17 +39,17 @@
// COMMON-LABEL: define{{.*}} spir_kernel void @device_side_enqueue(i32 addrspace(1)* align 4 %{{.*}}, i32 addrspace(1)* align 4 %b, i32 %i)
kernel void device_side_enqueue(global int *a, global int *b, int i) {
- // COMMON: %default_queue = alloca %opencl.queue_t*
+ // COMMON: %default_queue = alloca target("spirv.Queue")
queue_t default_queue;
// COMMON: %flags = alloca i32
unsigned flags = 0;
// COMMON: %ndrange = alloca %struct.ndrange_t
ndrange_t ndrange;
- // COMMON: %clk_event = alloca %opencl.clk_event_t*
+ // COMMON: %clk_event = alloca target("spirv.DeviceEvent")
clk_event_t clk_event;
- // COMMON: %event_wait_list = alloca %opencl.clk_event_t*
+ // COMMON: %event_wait_list = alloca target("spirv.DeviceEvent")
clk_event_t event_wait_list;
- // COMMON: %event_wait_list2 = alloca [1 x %opencl.clk_event_t*]
+ // COMMON: %event_wait_list2 = alloca [1 x target("spirv.DeviceEvent")]
clk_event_t event_wait_list2[] = {clk_event};
// COMMON: [[NDR:%[a-z0-9]+]] = alloca %struct.ndrange_t, align 4
@@ -77,14 +77,14 @@
// CHECK-LIFETIMES: %[[BLOCK_SIZES7:.*]] = alloca [1 x i64]
// Emits block literal on stack and block kernel [[INVLK1]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: store i8 addrspace(4)* addrspacecast (i8* bitcast (void (i8 addrspace(4)*)* [[INVL1:@__device_side_enqueue_block_invoke[^ ]*]] to i8*) to i8 addrspace(4)*), i8 addrspace(4)** %block.invoke
// B32: [[BL:%[0-9]+]] = bitcast <{ i32, i32, i8 addrspace(4)*, i32 addrspace(1)*, i32, i32 addrspace(1)* }>* %block to %struct.__opencl_block_literal_generic*
// B64: [[BL:%[0-9]+]] = bitcast <{ i32, i32, i8 addrspace(4)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 }>* %block to %struct.__opencl_block_literal_generic*
// COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast %struct.__opencl_block_literal_generic* [[BL]] to i8 addrspace(4)*
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
- // COMMON-SAME: %opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVLK1:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange,
@@ -93,15 +93,15 @@
});
// Emits block literal on stack and block kernel [[INVLK2]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
- // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast %opencl.clk_event_t{{.*}}** %event_wait_list to %opencl.clk_event_t{{.*}}* addrspace(4)*
- // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)*
+ // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast target("spirv.DeviceEvent")* %event_wait_list to target("spirv.DeviceEvent") addrspace(4)*
+ // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast target("spirv.DeviceEvent")* %clk_event to target("spirv.DeviceEvent") addrspace(4)*
// COMMON: store i8 addrspace(4)* addrspacecast (i8* bitcast (void (i8 addrspace(4)*)* [[INVL2:@__device_side_enqueue_block_invoke[^ ]*]] to i8*) to i8 addrspace(4)*), i8 addrspace(4)** %block.invoke
// COMMON: [[BL:%[0-9]+]] = bitcast <{ i32, i32, i8 addrspace(4)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block4 to %struct.__opencl_block_literal_generic*
// COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast %struct.__opencl_block_literal_generic* [[BL]] to i8 addrspace(4)*
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
- // COMMON-SAME: (%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]],
+ // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, target("spirv.DeviceEvent") addrspace(4)* [[WAIT_EVNT]], target("spirv.DeviceEvent") addrspace(4)* [[EVNT]],
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVLK2:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
@@ -110,14 +110,14 @@
});
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
- // COMMON-SAME: (%opencl.queue_t{{.*}}* {{%[0-9]+}}, i32 {{%[0-9]+}}, %struct.ndrange_t* {{.*}}, i32 1, %opencl.clk_event_t{{.*}}* addrspace(4)* null, %opencl.clk_event_t{{.*}}* addrspace(4)* null,
+ // COMMON-SAME: (target("spirv.Queue") {{%[0-9]+}}, i32 {{%[0-9]+}}, %struct.ndrange_t* {{.*}}, i32 1, target("spirv.DeviceEvent") addrspace(4)* null, target("spirv.DeviceEvent") 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 %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [1 x i64]* %[[BLOCK_SIZES1]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull [[LIFETIME_PTR]])
@@ -129,7 +129,7 @@
// B64: %[[TMP:.*]] = getelementptr [1 x i64], [1 x i64]* %[[BLOCK_SIZES1]], i32 0, i32 0
// B64: store i64 256, i64* %[[TMP]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
- // COMMON-SAME: %opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK1:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG1]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1,
// B32-SAME: i32* %[[TMP]])
@@ -142,7 +142,7 @@
char c;
// Emits global block literal [[BLG2]] and block kernel [[INVGK2]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [1 x i64]* %[[BLOCK_SIZES2]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull [[LIFETIME_PTR]])
@@ -154,7 +154,7 @@
// B64: %[[TMP:.*]] = getelementptr [1 x i64], [1 x i64]* %[[BLOCK_SIZES2]], i32 0, i32 0
// B64: store i64 %{{.*}}, i64* %[[TMP]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
- // COMMON-SAME: %opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK2:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG2]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1,
// B32-SAME: i32* %[[TMP]])
@@ -166,11 +166,11 @@
c);
// Emits global block literal [[BLG3]] and block kernel [[INVGK3]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
- // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
- // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast %opencl.clk_event_t{{.*}}** [[AD]] to %opencl.clk_event_t{{.*}}* addrspace(4)*
- // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)*
+ // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], [1 x target("spirv.DeviceEvent")]* %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+ // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast target("spirv.DeviceEvent")* [[AD]] to target("spirv.DeviceEvent") addrspace(4)*
+ // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast target("spirv.DeviceEvent")* %clk_event to target("spirv.DeviceEvent") addrspace(4)*
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [1 x i64]* %[[BLOCK_SIZES3]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull [[LIFETIME_PTR]])
// CHECK-LIFETIMES-NEXT: getelementptr inbounds [1 x i64], [1 x i64]* %[[BLOCK_SIZES3]], i64 0, i64 0
@@ -181,7 +181,7 @@
// B64: %[[TMP:.*]] = getelementptr [1 x i64], [1 x i64]* %[[BLOCK_SIZES3]], i32 0, i32 0
// B64: store i64 256, i64* %[[TMP]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
- // COMMON-SAME: (%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]],
+ // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, target("spirv.DeviceEvent") addrspace(4)* [[WAIT_EVNT]], target("spirv.DeviceEvent") addrspace(4)* [[EVNT]],
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK3:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG3]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1,
// B32-SAME: i32* %[[TMP]])
@@ -193,11 +193,11 @@
256);
// Emits global block literal [[BLG4]] and block kernel [[INVGK4]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
- // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
- // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast %opencl.clk_event_t{{.*}}** [[AD]] to %opencl.clk_event_t{{.*}}* addrspace(4)*
- // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)*
+ // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], [1 x target("spirv.DeviceEvent")]* %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+ // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast target("spirv.DeviceEvent")* [[AD]] to target("spirv.DeviceEvent") addrspace(4)*
+ // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast target("spirv.DeviceEvent")* %clk_event to target("spirv.DeviceEvent") addrspace(4)*
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [1 x i64]* %[[BLOCK_SIZES4]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull [[LIFETIME_PTR]])
// CHECK-LIFETIMES-NEXT: getelementptr inbounds [1 x i64], [1 x i64]* %[[BLOCK_SIZES4]], i64 0, i64 0
@@ -208,7 +208,7 @@
// B64: %[[TMP:.*]] = getelementptr [1 x i64], [1 x i64]* %[[BLOCK_SIZES4]], i32 0, i32 0
// B64: store i64 %{{.*}}, i64* %[[TMP]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
- // COMMON-SAME: (%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]],
+ // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* {{.*}}, i32 2, target("spirv.DeviceEvent") addrspace(4)* [[WAIT_EVNT]], target("spirv.DeviceEvent") addrspace(4)* [[EVNT]],
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK4:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG4]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1,
// B32-SAME: i32* %[[TMP]])
@@ -221,7 +221,7 @@
long l;
// Emits global block literal [[BLG5]] and block kernel [[INVGK5]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [1 x i64]* %[[BLOCK_SIZES5]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull [[LIFETIME_PTR]])
@@ -233,7 +233,7 @@
// B64: %[[TMP:.*]] = getelementptr [1 x i64], [1 x i64]* %[[BLOCK_SIZES5]], i32 0, i32 0
// B64: store i64 %{{.*}}, i64* %[[TMP]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
- // COMMON-SAME: (%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK5:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG5]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1,
// B32-SAME: i32* %[[TMP]])
@@ -245,7 +245,7 @@
l);
// Emits global block literal [[BLG6]] and block kernel [[INVGK6]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [3 x i64]* %[[BLOCK_SIZES6]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull [[LIFETIME_PTR]])
@@ -265,7 +265,7 @@
// B64: %[[BLOCK_SIZES63:.*]] = getelementptr [3 x i64], [3 x i64]* %[[BLOCK_SIZES6]], i32 0, i32 2
// B64: store i64 4, i64* %[[BLOCK_SIZES63]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
- // COMMON-SAME: (%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK6:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG6]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 3,
// B32-SAME: i32* %[[TMP]])
@@ -277,7 +277,7 @@
1, 2, 4);
// Emits global block literal [[BLG7]] and block kernel [[INVGK7]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// CHECK-LIFETIMES: [[LIFETIME_PTR:%[0-9]+]] = bitcast [1 x i64]* %[[BLOCK_SIZES7]] to i8*
// CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull [[LIFETIME_PTR]])
@@ -289,7 +289,7 @@
// B64: %[[TMP:.*]] = getelementptr [1 x i64], [1 x i64]* %[[BLOCK_SIZES7]], i32 0, i32 0
// B64: store i64 4294967296, i64* %[[TMP]], align 8
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
- // COMMON-SAME: (%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK7:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG7]] to i8 addrspace(1)*) to i8 addrspace(4)*), i32 1,
// B32-SAME: i32* %[[TMP]])
@@ -319,10 +319,10 @@
block_A();
// Emits global block literal [[BLG8]] and block kernel [[INVGK8]]. [[INVGK8]] calls [[INVG8]].
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
- // COMMON-SAME: %opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVGK8:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* [[BLG8]] to i8 addrspace(1)*) to i8 addrspace(4)*))
enqueue_kernel(default_queue, flags, ndrange, block_A);
@@ -365,11 +365,11 @@
};
// Emits block literal on stack and block kernel [[INVLK3]].
// COMMON: store i8 addrspace(4)* addrspacecast (i8* bitcast (void (i8 addrspace(4)*)* [[INVL3:@__device_side_enqueue_block_invoke[^ ]*]] to i8*) to i8 addrspace(4)*), i8 addrspace(4)** %block.invoke
- // COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
+ // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), target("spirv.Queue")* %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast %struct.__opencl_block_literal_generic* {{.*}} to i8 addrspace(4)*
// COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
- // COMMON-SAME: %opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+ // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], %struct.ndrange_t* byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
// COMMON-SAME: i8 addrspace(4)* addrspacecast (i8* bitcast ({{.*}} [[INVLK3:[^ ]+_kernel]] to i8*) to i8 addrspace(4)*),
// COMMON-SAME: i8 addrspace(4)* [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange, block_C);
Index: clang/test/CodeGenOpenCL/cast_image.cl
===================================================================
--- clang/test/CodeGenOpenCL/cast_image.cl
+++ 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
+// RUNx: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s
#ifdef __AMDGCN__
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -234,7 +234,12 @@
}
void CodeGenModule::createOpenCLRuntime() {
- OpenCLRuntime.reset(new CGOpenCLRuntime(*this));
+ // Use a different type mapping scheme for when SPIR-V wants to use target
+ // extension types.
+ if (getTriple().isSPIRV() || getTriple().isSPIR())
+ OpenCLRuntime.reset(new CGSpirVOpenCLRuntime(*this));
+ else
+ OpenCLRuntime.reset(new CGOpenCLRuntime(*this));
}
void CodeGenModule::createOpenMPRuntime() {
Index: clang/lib/CodeGen/CGOpenCLRuntime.h
===================================================================
--- clang/lib/CodeGen/CGOpenCLRuntime.h
+++ clang/lib/CodeGen/CGOpenCLRuntime.h
@@ -38,7 +38,7 @@
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 @@
virtual llvm::Type *getPipeType(const PipeType *T);
- llvm::PointerType *getSamplerType(const Type *T);
+ virtual llvm::Type *getSamplerType(const Type *T);
// Returns a value which indicates the size in bytes of the pipe
// element.
@@ -101,6 +101,23 @@
llvm::Function *getInvokeFunction(const Expr *E);
};
+class CGSpirVOpenCLRuntime : public CGOpenCLRuntime {
+protected:
+ virtual llvm::Type *getPipeType(const PipeType *T, StringRef Name,
+ llvm::Type *&PipeTy) override {
+ return CGOpenCLRuntime::getPipeType(T, Name, PipeTy);
+ }
+
+public:
+ CGSpirVOpenCLRuntime(CodeGenModule &CGM) : CGOpenCLRuntime(CGM) {}
+ virtual ~CGSpirVOpenCLRuntime();
+
+ virtual llvm::Type *convertOpenCLSpecificType(const Type *T) override;
+
+ virtual llvm::Type *getPipeType(const PipeType *T) override;
+
+ virtual llvm::Type *getSamplerType(const Type *T) override;
+};
}
}
Index: clang/lib/CodeGen/CGOpenCLRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenCLRuntime.cpp
+++ clang/lib/CodeGen/CGOpenCLRuntime.cpp
@@ -31,8 +31,7 @@
}
llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
- assert(T->isOpenCLSpecificType() &&
- "Not an OpenCL specific type!");
+ assert(T->isOpenCLSpecificType() && "Not an OpenCL specific type!");
switch (cast<BuiltinType>(T)->getKind()) {
default:
@@ -91,12 +90,13 @@
return PipeTy;
}
-llvm::PointerType *CGOpenCLRuntime::getSamplerType(const Type *T) {
- if (!SamplerTy)
+llvm::Type *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)));
+ CGM.getContext().getOpenCLTypeAddrSpace(T)));
+ }
return SamplerTy;
}
@@ -189,3 +189,80 @@
EnqueuedBlockMap[Block].Kernel = F;
return EnqueuedBlockMap[Block];
}
+
+CGSpirVOpenCLRuntime::~CGSpirVOpenCLRuntime() {}
+
+/// Construct a SPIR-V target extension type for the given OpenCL image type.
+static llvm::Type *getSPIRVType(llvm::LLVMContext &Ctx, StringRef BaseType,
+ StringRef OpenCLName, unsigned AccessQualifier) {
+ SmallVector<unsigned, 7> IntParams = {0, 0, 0, 0, 0, 0};
+
+ // Choose the dimension of the image--this corresponds to the Dim parameter,
+ // so (e.g.) a 2D image has value 1, not 2.
+ if (OpenCLName.startswith("image2d"))
+ IntParams[0] = 1;
+ else if (OpenCLName.startswith("image3d"))
+ IntParams[0] = 2;
+ else if (OpenCLName == "image1d_buffer")
+ IntParams[0] = 5;
+ else
+ assert(OpenCLName.startswith("image1d") && "Unknown image type");
+
+ // Other boolean 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 *CGSpirVOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
+ assert(T->isOpenCLSpecificType() && "Not an OpenCL specific type!");
+
+ llvm::LLVMContext &Ctx = CGM.getLLVMContext();
+ enum AccessQualifier : unsigned { AQ_ro = 0, AQ_wo = 1, AQ_rw = 2 };
+ switch (cast<BuiltinType>(T)->getKind()) {
+ default:
+ llvm_unreachable("Unexpected opencl builtin type!");
+ return nullptr;
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
+ case BuiltinType::Id: \
+ return getSPIRVType(Ctx, "spirv.Image", #ImgType, AQ_##Suffix);
+#include "clang/Basic/OpenCLImageTypes.def"
+ case BuiltinType::OCLSampler:
+ return getSamplerType(T);
+ 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"
+ }
+}
+
+llvm::Type *CGSpirVOpenCLRuntime::getPipeType(const PipeType *T) {
+ llvm::Type *&TargetTy = T->isReadOnly() ? PipeROTy : PipeWOTy;
+ return TargetTy = llvm::TargetExtType::get(CGM.getLLVMContext(), "spirv.Pipe",
+ {}, {!T->isReadOnly()});
+}
+
+llvm::Type *CGSpirVOpenCLRuntime::getSamplerType(const Type *T) {
+ if (!SamplerTy) {
+ SamplerTy = llvm::TargetExtType::get(CGM.getLLVMContext(), "spirv.Sampler");
+ }
+ return SamplerTy;
+}
+
Index: clang/lib/CodeGen/CGExprScalar.cpp
===================================================================
--- clang/lib/CodeGen/CGExprScalar.cpp
+++ clang/lib/CodeGen/CGExprScalar.cpp
@@ -2253,12 +2253,18 @@
case CK_FunctionToPointerDecay:
return EmitLValue(E).getPointer(CGF);
- case CK_NullToPointer:
+ case CK_NullToPointer: {
if (MustVisitNullValue(E))
CGF.EmitIgnoredExpr(E);
- return CGF.CGM.getNullPointer(cast<llvm::PointerType>(ConvertType(DestTy)),
- DestTy);
+ // The type may be a target extension type instead of a pointer type
+ // (e.g., OpenCL types mapped for SPIR-V). In the former case, emit a
+ // null value instead.
+ llvm::Type *LlvmTy = ConvertType(DestTy);
+ if (auto *PointerTy = dyn_cast<llvm::PointerType>(LlvmTy))
+ return CGF.CGM.getNullPointer(PointerTy, DestTy);
+ return llvm::Constant::getNullValue(LlvmTy);
+ }
case CK_NullToMemberPointer: {
if (MustVisitNullValue(E))
Index: clang/include/clang/Basic/OpenCLExtensionTypes.def
===================================================================
--- clang/include/clang/Basic/OpenCLExtensionTypes.def
+++ clang/include/clang/Basic/OpenCLExtensionTypes.def
@@ -28,10 +28,10 @@
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
Index: clang/include/clang-c/Index.h
===================================================================
--- clang/include/clang-c/Index.h
+++ clang/include/clang-c/Index.h
@@ -34,7 +34,7 @@
* compatible, thus CINDEX_VERSION_MAJOR is expected to remain stable.
*/
#define CINDEX_VERSION_MAJOR 0
-#define CINDEX_VERSION_MINOR 62
+#define CINDEX_VERSION_MINOR 63
#define CINDEX_VERSION_ENCODE(major, minor) (((major)*10000) + ((minor)*1))
@@ -2781,10 +2781,15 @@
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,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
