[CIR][OpenCL] Support lowering of OCL opaque types

clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h

@@ -155,9 +155,10 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
       return cir::ZeroAttr::get(RecordTy);
     if (auto methodTy = mlir::dyn_cast<cir::MethodType>(ty))
       return getNullMethodAttr(methodTy);
-    if (mlir::isa<cir::BoolType>(ty)) {
+    if (mlir::isa<cir::BoolType>(ty))
       return getFalseAttr();
-    }
+    if (mlir::isa<cir::OpaqueType>(ty))
+      return cir::ZeroAttr::get(ty);
     llvm_unreachable("Zero initializer for given type is NYI");
   }
 

clang/include/clang/CIR/Dialect/IR/CIRTypes.td

@@ -751,6 +751,48 @@ def CIR_RecordType : CIR_Type<"Record", "record", [
 def CIRRecordType : Type<
   CPred<"::mlir::isa<::cir::RecordType>($_self)">, "CIR record type">;
 
+//===----------------------------------------------------------------------===//
+// Minimal opaque type (used for OpenCL opaque builtin types)
+//===----------------------------------------------------------------------===//
+
+def CIR_OCLOpaqueType : CIR_Type<"Opaque", "opaque"> {
+  let summary = "Named opaque type for OpenCL-style builtin opaque objects";
+
+  let description = [{
+    Represents a target-independent opaque type used for OpenCL opaque
+    builtin types such as `event_t`, `sampler_t`, `clk_event_t` and `queue_t`.
+
+    The type has no defined size or layout. CIR carries it through
+    lowering and delegates the final representation to the target codegen
+    (e.g. SPIR/SPIR-V lowering), which maps the logical opaque kind to
+    the correct LLVM type.
+
+    The `tag` attribute identifies the opaque category (e.g. `"event"`).
+    Values of this type typically appear only through pointer types.
+
+    Example:
+      !cir.ptr<!cir.opaque<"event">, addrspace(1)>
+  }];
+
+  let parameters = (ins "mlir::StringAttr":$tag);
+
+  let builders = [
+    TypeBuilder<(ins "mlir::StringAttr":$tag), [{
+      return $_get($_ctxt, tag);
+    }]>
+  ];
+
+  let extraClassDeclaration = [{
+    static llvm::StringRef getEventTag() { return "event"; }
+  }];
+
+  let assemblyFormat = [{
+    `<` $tag `>`
+  }];
+
+  let skipDefaultBuilders = 1;
+}
+
 //===----------------------------------------------------------------------===//
 // Global type constraints
 //===----------------------------------------------------------------------===//
@@ -759,7 +801,7 @@ def CIR_AnyType : AnyTypeOf<[
   CIR_IntType, CIR_PointerType, CIR_DataMemberType, CIR_MethodType,
   CIR_BoolType, CIR_ArrayType, CIR_VectorType, CIR_FuncType, CIR_VoidType,
   CIR_RecordType, CIR_ExceptionType, CIR_AnyFloatType, CIR_ComplexType,
-  CIR_VPtrType
+  CIR_VPtrType, CIR_OCLOpaqueType
 ]>;
 
 #endif // MLIR_CIR_DIALECT_CIR_TYPES

clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp

@@ -1946,7 +1946,10 @@ mlir::Value ScalarExprEmitter::VisitCastExpr(CastExpr *CE) {
                                          DestTy);
   }
   case CK_ZeroToOCLOpaqueType:
-    llvm_unreachable("NYI");
+    // OpenCL: event_t e = async_work_group_copy(..., 0);
+    // The source is an integer constant zero; the destination is an OpenCL
+    // opaque type
+    return emitNullValue(DestTy, CGF.getLoc(E->getExprLoc()));
   case CK_IntToOCLSampler:
     llvm_unreachable("NYI");
 

clang/lib/CIR/CodeGen/CIRGenTypes.cpp

@@ -512,13 +512,23 @@ mlir::Type CIRGenTypes::convertType(QualType T) {
 #include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
-    case BuiltinType::OCLSampler:
     case BuiltinType::OCLEvent:
+      ResultType = cir::OpaqueType::get(
+          Builder.getContext(),
+          mlir::StringAttr::get(Builder.getContext(),
+                                cir::OpaqueType::getEventTag()));
+      break;
+    case BuiltinType::OCLSampler:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
-    case BuiltinType::OCLReserveID:
-      assert(0 && "not implemented");
+      llvm_unreachable("NYI");
       break;
+    case BuiltinType::OCLReserveID:
+      ResultType = cir::RecordType::get(
+          &getMLIRContext(), {},
+          mlir::StringAttr::get(&getMLIRContext(), "ocl_reserve_id"), false,
+          false, cir::RecordType::Struct);
+
     case BuiltinType::SveInt8:
     case BuiltinType::SveUint8:
     case BuiltinType::SveInt8x2:

clang/lib/CIR/Dialect/IR/CIRDialect.cpp

@@ -392,7 +392,7 @@ static LogicalResult checkConstantTypes(mlir::Operation *op, mlir::Type opType,
 
   if (isa<cir::ZeroAttr>(attrType)) {
     if (::mlir::isa<cir::RecordType, cir::ArrayType, cir::ComplexType,
-                    cir::VectorType>(opType))
+                    cir::VectorType, cir::OpaqueType>(opType))
       return success();
     return op->emitOpError("zero expects record or array type");
   }

clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h

@@ -34,6 +34,10 @@ class TargetLoweringInfo {
 
   virtual unsigned
   getTargetAddrSpaceFromCIRAddrSpace(cir::AddressSpace addrSpace) const = 0;
+
+  virtual mlir::Type getOpaqueType(cir::OpaqueType type) const {
+    llvm_unreachable("NYI");
+  }
 };
 
 } // namespace cir

clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp

@@ -16,6 +16,7 @@
 #include "clang/CIR/MissingFeatures.h"
 #include "llvm/Support/Casting.h"
 #include "llvm/Support/ErrorHandling.h"
+#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
 
 using ABIArgInfo = cir::ABIArgInfo;
 using MissingFeature = cir::MissingFeatures;
@@ -60,6 +61,11 @@ class AMDGPUTargetLoweringInfo : public TargetLoweringInfo {
       cir_cconv_unreachable("Unknown CIR address space for this target");
     }
   }
+
+  mlir::Type getOpaqueType(cir::OpaqueType type) const override {
+    assert(!cir::MissingFeatures::addressSpace());
+    return mlir::LLVM::LLVMPointerType::get(type.getContext());
+  }
 };
 
 } // namespace

clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp

@@ -12,8 +12,10 @@
 #include "TargetInfo.h"
 #include "TargetLoweringInfo.h"
 #include "clang/CIR/ABIArgInfo.h"
+#include "clang/CIR/Dialect/IR/CIRTypes.h"
 #include "clang/CIR/MissingFeatures.h"
 #include "llvm/Support/ErrorHandling.h"
+#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
 
 using ABIArgInfo = cir::ABIArgInfo;
 using MissingFeature = cir::MissingFeatures;
@@ -58,6 +60,16 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
       cir_cconv_unreachable("Unknown CIR address space for this target");
     }
   }
+
+  mlir::Type getOpaqueType(cir::OpaqueType type) const override {
+    if (type.getTag() != cir::OpaqueType::getEventTag())
+      llvm_unreachable("NYI");
+
+    return mlir::LLVM::LLVMTargetExtType::get(type.getContext(),
+                                              /*extTypeName=*/"spirv.Event",
+                                              /*typeParams=*/{},
+                                              /*intParams=*/{});
+  }
 };
 
 } // namespace

clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp

@@ -2138,6 +2138,45 @@ mlir::LogicalResult CIRToLLVMConstantOpLowering::matchAndRewrite(
     rewriter.replaceOp(op, lowerCirAttrAsValue(op, op.getValue(), rewriter,
                                                getTypeConverter(), dataLayout));
     return mlir::success();
+  } else if (mlir::isa<cir::OpaqueType>(op.getType())) {
+    mlir::Attribute valAttr = op.getValue();
+    mlir::Type llvmTy = getTypeConverter()->convertType(op.getType());
+    // If the attribute is ZeroAttr or UndefAttr, handle it:
+    if (mlir::isa<cir::ZeroAttr, cir::UndefAttr>(valAttr)) {
+      // Handle target-ext type
+      if (auto tgtExtTy =
+              llvm::dyn_cast_or_null<mlir::LLVM::LLVMTargetExtType>(llvmTy)) {
+        // Produce a real zero constant if the target-ext type allows it
+        if (tgtExtTy.hasProperty(mlir::LLVM::LLVMTargetExtType::HasZeroInit)) {
+          if (mlir::isa<cir::ZeroAttr>(valAttr)) {
+            auto zero =
+                mlir::LLVM::ZeroOp::create(rewriter, op.getLoc(), llvmTy);
+            rewriter.replaceOp(op, zero.getResult());
+            return mlir::success();
+          }
+          // Fallback: emit an undef of that exact llvm type so users have
+          // matching types.
+          auto undef =
+              mlir::LLVM::UndefOp::create(rewriter, op.getLoc(), llvmTy);
+          rewriter.replaceOp(op, undef.getResult());
+          return mlir::success();
+        }
+      } else {
+        // Target ext type does not support zero init — use `ptr null` of
+        // the target-ext type (so users still have the expected type).
+        auto ptrTy = mlir::LLVM::LLVMPointerType::get(getContext());
+        auto nullPtr = mlir::LLVM::ZeroOp::create(rewriter, op.getLoc(), ptrTy);
+
+        rewriter.replaceOp(op, nullPtr.getResult());
+        return mlir::success();
+      }
+    }
+
+    // If the attr is a non-zero concrete value, we must decide if the target
+    // expects an encoded representation. Most target-ext types for OpenCL
+    // do not accept arbitrary non-zero constants; reject them.
+    return op.emitError() << "non-zero constant for target extension type "
+                          << llvmTy << " is unsupported";
   } else
     return op.emitError() << "unsupported constant type " << op.getType();
 
@@ -5138,6 +5177,10 @@ void prepareTypeConverter(mlir::LLVMTypeConverter &converter,
   converter.addConversion([&](cir::VoidType type) -> mlir::Type {
     return mlir::LLVM::LLVMVoidType::get(type.getContext());
   });
+
+  converter.addConversion([&](cir::OpaqueType type) -> mlir::Type {
+    return lowerModule->getTargetLoweringInfo().getOpaqueType(type);
+  });
 }
 
 void buildCtorDtorList(

clang/lib/CIR/Lowering/ThroughMLIR/LowerCIRToMLIR.cpp

@@ -1767,6 +1767,9 @@ static mlir::TypeConverter prepareTypeConverter() {
       return nullptr;
     return mlir::VectorType::get(2, elemTy);
   });
+  converter.addConversion([&](cir::OpaqueType type) -> mlir::Type {
+    llvm_unreachable("NYI");
+  });
   return converter;
 }
 

clang/test/CIR/CodeGen/OpenCL/async_copy.cl

@@ -0,0 +1,34 @@
+// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-cir -o - %s -fclangir | FileCheck %s --check-prefix=CIR-SPIR
+// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s -fclangir | FileCheck %s --check-prefix=LLVM-SPIR
+// RUN: %clang -cc1 -triple spirv64-unknown-unknown -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s | FileCheck %s --check-prefix=OG-LLVM-SPIR
+
+// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -finclude-default-header -emit-cir -o - %s -fclangir | FileCheck %s --check-prefix=CIR-AMDGCN
+// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s -fclangir | FileCheck %s --check-prefix=LLVM-AMDGCN
+// RUN: %clang -cc1 -triple amdgcn-amd-amdhsa -cl-std=CL2.0 -finclude-default-header -emit-llvm -o - %s | FileCheck %s --check-prefix=OG-LLVM-AMDGCN
+
+
+// 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-SPIR: cir.call @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!cir.ptr<!s32i, addrspace(offload_local)>, !cir.ptr<!s32i, addrspace(offload_global)>, !u64i, !cir.opaque<"event">) -> !cir.opaque<"event">
+// LLVM-SPIR: call spir_func target("spirv.Event") @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i64 %{{.*}}, target("spirv.Event") zeroinitializer)
+// OG-LLVM-SPIR: 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
+
+// CIR-AMDGCN: cir.call @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!cir.ptr<!s32i, addrspace(offload_local)>, !cir.ptr<!s32i, addrspace(offload_global)>, !u64i, !cir.opaque<"event">) -> !cir.opaque<"event">
+// LLVM-AMDGCN: call ptr @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i64 %{{.*}}, ptr null)
+// OG-LLVM-AMDGCN: call ptr @_Z21async_work_group_copyPU3AS3iPU3AS1Kim9ocl_event(ptr addrspace(3) noundef %{{.*}}, ptr addrspace(1) noundef %{{.*}}, i64 noundef %{{.*}}, ptr null)