jcranmer-intel updated this revision to Diff 495866. jcranmer-intel marked 2 inline comments as done. jcranmer-intel added a comment. Herald added a subscriber: jdoerfert.
This updates code, and rebases tests on top of trunk. Note: test issues still haven't been fixed, will fix that likely in the next revision of the patch. 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/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
Index: llvm/docs/SPIRVUsage.rst =================================================================== --- llvm/docs/SPIRVUsage.rst +++ llvm/docs/SPIRVUsage.rst @@ -75,3 +75,36 @@ 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)``. 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 @@ -38,17 +38,17 @@ // 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 + // 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 ptr + // COMMON: %clk_event = alloca target("spirv.DeviceEvent") clk_event_t clk_event; - // COMMON: %event_wait_list = alloca ptr + // COMMON: %event_wait_list = alloca target("spirv.DeviceEvent") clk_event_t event_wait_list; - // COMMON: %event_wait_list2 = alloca [1 x ptr] + // 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 @@ -76,12 +76,12 @@ // 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 + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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-SAME: target("spirv.Queue") [[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 +90,14 @@ }); // Emits block literal on stack and block kernel [[INVLK2]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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-SAME: (target("spirv.Queue") [[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, @@ -106,14 +106,14 @@ }); // 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-SAME: (target("spirv.Queue") {{%[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 + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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( @@ -123,7 +123,7 @@ // 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-SAME: target("spirv.Queue") [[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,7 +136,7 @@ char c; // Emits global block literal [[BLG2]] and block kernel [[INVGK2]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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( @@ -146,7 +146,7 @@ // 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-SAME: target("spirv.Queue") [[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,9 +158,9 @@ c); // Emits global block literal [[BLG3]] and block kernel [[INVGK3]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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 + // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], 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]]) @@ -171,7 +171,7 @@ // 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-SAME: (target("spirv.Queue") [[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,9 +183,9 @@ 256); // Emits global block literal [[BLG4]] and block kernel [[INVGK4]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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 + // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], 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]]) @@ -196,7 +196,7 @@ // 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-SAME: (target("spirv.Queue") [[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,7 +209,7 @@ long l; // Emits global block literal [[BLG5]] and block kernel [[INVGK5]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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( @@ -219,7 +219,7 @@ // 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-SAME: (target("spirv.Queue") [[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,7 +231,7 @@ l); // Emits global block literal [[BLG6]] and block kernel [[INVGK6]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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( @@ -249,7 +249,7 @@ // 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-SAME: (target("spirv.Queue") [[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,7 +261,7 @@ 1, 2, 4); // Emits global block literal [[BLG7]] and block kernel [[INVGK7]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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( @@ -271,7 +271,7 @@ // 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-SAME: (target("spirv.Queue") [[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]]) @@ -301,10 +301,10 @@ block_A(); // Emits global block literal [[BLG8]] and block kernel [[INVGK8]]. [[INVGK8]] calls [[INVG8]]. - // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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-SAME: target("spirv.Queue") [[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); @@ -347,11 +347,11 @@ }; // 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 + // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), 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-SAME: target("spirv.Queue") [[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); Index: clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl =================================================================== --- clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl +++ clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl @@ -27,7 +27,7 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 // // 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/TargetInfo.h =================================================================== --- clang/lib/CodeGen/TargetInfo.h +++ clang/lib/CodeGen/TargetInfo.h @@ -379,6 +379,12 @@ // DO NOTHING by default. return false; } + + /// Return an LLVM target extension type that corresponds to an OpenCL type, + /// if such a type is necessary. + virtual llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const { + return nullptr; + } }; } // namespace CodeGen Index: clang/lib/CodeGen/TargetInfo.cpp =================================================================== --- clang/lib/CodeGen/TargetInfo.cpp +++ clang/lib/CodeGen/TargetInfo.cpp @@ -10561,6 +10561,7 @@ } unsigned getOpenCLKernelCallingConv() const override; + llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const override; }; class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo { public: @@ -10975,6 +10976,76 @@ return false; } +/// 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) { + // These parameters compare to the operands of OpTypeImage (see + // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeImage + // for more details). + SmallVector<unsigned, 7> IntParams = {0, 0, 0, 0, 0, 0}; + + // Choose the dimension of the image--this corresponds to the Dim parameter + // (https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_dim), or + // the instances of Dim in llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td. + 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"); + + // 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 *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 getSPIRVType(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 //===----------------------------------------------------------------------===// 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); + llvm::Type *getSamplerType(const Type *T); // Returns a value which indicates the size in bytes of the pipe // element. Index: clang/lib/CodeGen/CGOpenCLRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenCLRuntime.cpp +++ clang/lib/CodeGen/CGOpenCLRuntime.cpp @@ -31,8 +31,11 @@ } 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::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,17 @@ 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) { + 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; } 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)) @@ -2787,10 +2787,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
