From dc0ffd70d8112431b6bcd218ba4d0f0c4318d26a Mon Sep 17 00:00:00 2001 From: Mahmood Yassin Date: Mon, 24 Nov 2025 15:31:44 +0200 Subject: [PATCH 1/2] [CIR][OpenCL] Support lowering of OCL opaque types Implement handling for zero-initialization casts to OpenCL opaque types in CIR. This covers cases like `event_t e = async_work_group_copy(..., 0)`. - `VisitCastExpr`: CK_ZeroToOCLOpaqueType now returns a null pointer of the appropriate opaque type instead of `llvm_unreachable`. - `CIRGenTypes::convertType`: Added proper CIR type conversions for OpenCL opaque types including event, queue, and reserve_id types. - Provides consistent CIR representation for OpenCL opaque objects. --- clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 14 ++++++++++++-- clang/lib/CIR/CodeGen/CIRGenTypes.cpp | 7 ++++++- 2 files changed, 18 insertions(+), 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index a1adea6cd172..f815a9871476 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp @@ -1945,8 +1945,18 @@ mlir::Value ScalarExprEmitter::VisitCastExpr(CastExpr *CE) { return emitComplexToScalarConversion(CGF.getLoc(CE->getExprLoc()), V, Kind, DestTy); } - case CK_ZeroToOCLOpaqueType: - llvm_unreachable("NYI"); + case CK_ZeroToOCLOpaqueType: { + // OpenCL: event_t e = async_work_group_copy(..., 0); + // The source is an integer constant zero; the destination is an OpenCL + // opaque type + mlir::Type destTy = CGF.convertType(DestTy); + auto PtrTy = + cir::PointerType::get(destTy, cir::AddressSpace::OffloadPrivate); + auto constNullPtrAttr = Builder.getConstNullPtrAttr(PtrTy); + auto nullVal = + Builder.getConstant(CGF.getLoc(E->getExprLoc()), constNullPtrAttr); + return nullVal; + } case CK_IntToOCLSampler: llvm_unreachable("NYI"); diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp index 4de89fc7081a..8d2068b4e9fd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp @@ -516,8 +516,13 @@ mlir::Type CIRGenTypes::convertType(QualType T) { case BuiltinType::OCLEvent: case BuiltinType::OCLClkEvent: case BuiltinType::OCLQueue: + ResultType = Builder.getVoidPtrTy(); + break; case BuiltinType::OCLReserveID: - assert(0 && "not implemented"); + ResultType = cir::RecordType::get( + &getMLIRContext(), {}, + mlir::StringAttr::get(&getMLIRContext(), "ocl_reserve_id"), false, + false, cir::RecordType::Struct); break; case BuiltinType::SveInt8: case BuiltinType::SveUint8: From 0eb95502663b90842cae7bfd8ff0e5ee1d108665 Mon Sep 17 00:00:00 2001 From: Mahmood Yassin Date: Tue, 25 Nov 2025 13:12:31 +0200 Subject: [PATCH 2/2] adding test --- clang/test/CIR/CodeGen/OpenCL/async_copy.cl | 25 +++++++++++++++++++++ 1 file changed, 25 insertions(+) create mode 100644 clang/test/CIR/CodeGen/OpenCL/async_copy.cl diff --git a/clang/test/CIR/CodeGen/OpenCL/async_copy.cl b/clang/test/CIR/CodeGen/OpenCL/async_copy.cl new file mode 100644 index 000000000000..256d28f1fd3f --- /dev/null +++ b/clang/test/CIR/CodeGen/OpenCL/async_copy.cl @@ -0,0 +1,25 @@ +// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-cir -o - %s -fclangir | FileCheck %s --check-prefix=CIR +// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s -fclangir | FileCheck %s --check-prefix=LLVM +// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s | FileCheck %s --check-prefix=OG-LLVM + +// Simple kernel using async_work_group_copy + wait_group_events + +__kernel void test_async_copy(__global int *g_in, __local int *l_in, int size) { + // int gid = get_global_id(0); + + // Trigger async copy: global to local + // event_t e_in = + async_work_group_copy( + l_in, // local destination + g_in,// + gid * size, // global source + size, // number of elements + (event_t)0 // no dependency + ); + + // Wait for the async operation to complete + // wait_group_events(1, &e_in); +} + +// CIR: cir.call @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!cir.ptr, !cir.ptr, !u64i, !cir.ptr) -> !cir.ptr +// LLVM: call spir_func ptr @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i64 %{{.*}}, ptr null) +// OG-LLVM: call spir_func target("spirv.Event") @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}}, i64 noundef %{{.*}}, target("spirv.Event") zeroinitializer \ No newline at end of file