diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index a533bea6f1e6..6ed3733f88e4 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/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(ty)) return getNullMethodAttr(methodTy); - if (mlir::isa(ty)) { + if (mlir::isa(ty)) return getFalseAttr(); - } + if (mlir::isa(ty)) + return cir::ZeroAttr::get(ty); llvm_unreachable("Zero initializer for given type is NYI"); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index 33abc3ffc66f..881a43790606 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/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, 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 diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp index a1adea6cd172..aa06da12ef4f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp +++ b/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"); diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp index 4de89fc7081a..a2be3ce73c41 100644 --- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp +++ b/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: diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index 4387142ac8c5..cf693001f22a 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -392,7 +392,7 @@ static LogicalResult checkConstantTypes(mlir::Operation *op, mlir::Type opType, if (isa(attrType)) { if (::mlir::isa(opType)) + cir::VectorType, cir::OpaqueType>(opType)) return success(); return op->emitOpError("zero expects record or array type"); } diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h index 114d8cc0f697..99d6107d7d24 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetLoweringInfo.h +++ b/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 diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp index 7432972889ed..1f77284dee02 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/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 diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp index 0a4dc640decd..8528f52d7f8c 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIR.cpp +++ b/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 diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index ebe7877b303f..84edc1d89a00 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/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(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(valAttr)) { + // Handle target-ext type + if (auto tgtExtTy = + llvm::dyn_cast_or_null(llvmTy)) { + // Produce a real zero constant if the target-ext type allows it + if (tgtExtTy.hasProperty(mlir::LLVM::LLVMTargetExtType::HasZeroInit)) { + if (mlir::isa(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( diff --git a/clang/lib/CIR/Lowering/ThroughMLIR/LowerCIRToMLIR.cpp b/clang/lib/CIR/Lowering/ThroughMLIR/LowerCIRToMLIR.cpp index 5b4ec9c5cf16..63775581a450 100644 --- a/clang/lib/CIR/Lowering/ThroughMLIR/LowerCIRToMLIR.cpp +++ b/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; } 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..53f39484559a --- /dev/null +++ b/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, !cir.ptr, !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, !cir.ptr, !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) \ No newline at end of file