Index: include/clang/AST/OperationKinds.def =================================================================== --- include/clang/AST/OperationKinds.def +++ include/clang/AST/OperationKinds.def @@ -324,6 +324,8 @@ // Convert a pointer to a different address space. CAST_OPERATION(AddressSpaceConversion) +// Convert an integer initializer to an OpenCL sampler. +CAST_OPERATION(IntToOCLSampler) //===- Binary Operations -------------------------------------------------===// // Operators listed in order of precedence. Index: include/clang/Basic/DiagnosticGroups.td =================================================================== --- include/clang/Basic/DiagnosticGroups.td +++ include/clang/Basic/DiagnosticGroups.td @@ -871,3 +871,7 @@ def OptionIgnored : DiagGroup<"option-ignored">; def UnknownArgument : DiagGroup<"unknown-argument">; + +// A warning group for warnings about code that clang accepts when +// compiling OpenCL C/C++ but which is not compatible with the SPIR spec. +def SpirCompat : DiagGroup<"spir-compat">; \ No newline at end of file Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -7862,6 +7862,12 @@ "the event_t type can only be used with __private address space qualifier">; def err_expected_kernel_void_return_type : Error< "kernel must have void return type">; +def err_sampler_initializer_not_integer : Error< + "sampler_t initialization requires 32-bit integer, not %0">; +def warn_sampler_initializer_invalid_bits : Warning< + "sampler initializer has invalid %0 bits">, InGroup, DefaultIgnore; +def err_sampler_initializer_not_constant : Error< + "sampler_t initialization requires compile time constant">; def err_sampler_argument_required : Error< "sampler_t variable required - got %0">; def err_wrong_sampler_addressspace: Error< Index: lib/AST/ASTContext.cpp =================================================================== --- lib/AST/ASTContext.cpp +++ lib/AST/ASTContext.cpp @@ -1661,11 +1661,12 @@ Width = Target->getPointerWidth(0); Align = Target->getPointerAlign(0); break; - case BuiltinType::OCLSampler: - // Samplers are modeled as integers. - Width = Target->getIntWidth(); - Align = Target->getIntAlign(); + case BuiltinType::OCLSampler: { + auto AS = getTargetAddressSpace(LangAS::opencl_constant); + Width = Target->getPointerWidth(AS); + Align = Target->getPointerAlign(AS); break; + } case BuiltinType::OCLEvent: case BuiltinType::OCLClkEvent: case BuiltinType::OCLQueue: Index: lib/AST/Expr.cpp =================================================================== --- lib/AST/Expr.cpp +++ lib/AST/Expr.cpp @@ -1570,6 +1570,7 @@ case CK_ARCReclaimReturnedObject: case CK_ARCExtendBlockObject: case CK_ZeroToOCLEvent: + case CK_IntToOCLSampler: assert(!getType()->isBooleanType() && "unheralded conversion to bool"); goto CheckNoBasePath; @@ -2748,7 +2749,8 @@ CE->getCastKind() == CK_ToUnion || CE->getCastKind() == CK_ConstructorConversion || CE->getCastKind() == CK_NonAtomicToAtomic || - CE->getCastKind() == CK_AtomicToNonAtomic) + CE->getCastKind() == CK_AtomicToNonAtomic || + CE->getCastKind() == CK_IntToOCLSampler) return CE->getSubExpr()->isConstantInitializer(Ctx, false, Culprit); break; Index: lib/AST/ExprConstant.cpp =================================================================== --- lib/AST/ExprConstant.cpp +++ lib/AST/ExprConstant.cpp @@ -8035,6 +8035,7 @@ case CK_ZeroToOCLEvent: case CK_NonAtomicToAtomic: case CK_AddressSpaceConversion: + case CK_IntToOCLSampler: llvm_unreachable("invalid cast kind for integral value"); case CK_BitCast: @@ -8526,6 +8527,7 @@ case CK_ZeroToOCLEvent: case CK_NonAtomicToAtomic: case CK_AddressSpaceConversion: + case CK_IntToOCLSampler: llvm_unreachable("invalid cast kind for complex value"); case CK_LValueToRValue: Index: lib/CodeGen/CGDebugInfo.h =================================================================== --- lib/CodeGen/CGDebugInfo.h +++ lib/CodeGen/CGDebugInfo.h @@ -67,6 +67,7 @@ #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ llvm::DIType *SingletonId = nullptr; #include "clang/Basic/OpenCLImageTypes.def" + llvm::DIType *OCLSamplerDITy = nullptr; llvm::DIType *OCLEventDITy = nullptr; llvm::DIType *OCLClkEventDITy = nullptr; llvm::DIType *OCLQueueDITy = nullptr; Index: lib/CodeGen/CGDebugInfo.cpp =================================================================== --- lib/CodeGen/CGDebugInfo.cpp +++ lib/CodeGen/CGDebugInfo.cpp @@ -476,9 +476,8 @@ SingletonId); #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: - return DBuilder.createBasicType( - "opencl_sampler_t", CGM.getContext().getTypeSize(BT), - CGM.getContext().getTypeAlign(BT), llvm::dwarf::DW_ATE_unsigned); + return getOrCreateStructPtrType("opencl_sampler_t", + OCLSamplerDITy); case BuiltinType::OCLEvent: return getOrCreateStructPtrType("opencl_event_t", OCLEventDITy); case BuiltinType::OCLClkEvent: Index: lib/CodeGen/CGExpr.cpp =================================================================== --- lib/CodeGen/CGExpr.cpp +++ lib/CodeGen/CGExpr.cpp @@ -3584,6 +3584,7 @@ case CK_ARCExtendBlockObject: case CK_CopyAndAutoreleaseBlockObject: case CK_AddressSpaceConversion: + case CK_IntToOCLSampler: return EmitUnsupportedLValue(E, "unexpected cast lvalue"); case CK_Dependent: Index: lib/CodeGen/CGExprAgg.cpp =================================================================== --- lib/CodeGen/CGExprAgg.cpp +++ lib/CodeGen/CGExprAgg.cpp @@ -750,6 +750,7 @@ case CK_BuiltinFnToFnPtr: case CK_ZeroToOCLEvent: case CK_AddressSpaceConversion: + case CK_IntToOCLSampler: llvm_unreachable("cast kind invalid for aggregate types"); } } Index: lib/CodeGen/CGExprComplex.cpp =================================================================== --- lib/CodeGen/CGExprComplex.cpp +++ lib/CodeGen/CGExprComplex.cpp @@ -484,6 +484,7 @@ case CK_BuiltinFnToFnPtr: case CK_ZeroToOCLEvent: case CK_AddressSpaceConversion: + case CK_IntToOCLSampler: llvm_unreachable("invalid cast kind for complex value"); case CK_FloatingRealToComplex: Index: lib/CodeGen/CGExprConstant.cpp =================================================================== --- lib/CodeGen/CGExprConstant.cpp +++ lib/CodeGen/CGExprConstant.cpp @@ -690,6 +690,9 @@ case CK_ConstructorConversion: return C; + case CK_IntToOCLSampler: + llvm_unreachable("global sampler variables are not generated"); + case CK_Dependent: llvm_unreachable("saw dependent cast!"); case CK_BuiltinFnToFnPtr: Index: lib/CodeGen/CGExprScalar.cpp =================================================================== --- lib/CodeGen/CGExprScalar.cpp +++ lib/CodeGen/CGExprScalar.cpp @@ -1573,7 +1573,10 @@ return llvm::Constant::getNullValue(ConvertType(DestTy)); } - } + case CK_IntToOCLSampler: + return CGF.CGM.createOpenCLIntToSamplerConversion(E, CGF); + + } // end of switch llvm_unreachable("unknown scalar cast"); } Index: lib/CodeGen/CGOpenCLRuntime.h =================================================================== --- lib/CodeGen/CGOpenCLRuntime.h +++ lib/CodeGen/CGOpenCLRuntime.h @@ -33,9 +33,11 @@ protected: CodeGenModule &CGM; llvm::Type *PipeTy; + llvm::PointerType *SamplerTy; public: - CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM), PipeTy(nullptr) {} + CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM), PipeTy(nullptr), + SamplerTy(nullptr) {} virtual ~CGOpenCLRuntime(); /// Emit the IR required for a work-group-local variable declaration, and add @@ -47,6 +49,8 @@ virtual llvm::Type *convertOpenCLSpecificType(const Type *T); virtual llvm::Type *getPipeType(); + + llvm::PointerType *getSamplerType(); }; } Index: lib/CodeGen/CGOpenCLRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenCLRuntime.cpp +++ lib/CodeGen/CGOpenCLRuntime.cpp @@ -47,7 +47,7 @@ ImgAddrSpc); #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: - return llvm::IntegerType::get(Ctx, 32); + return getSamplerType(); case BuiltinType::OCLEvent: return llvm::PointerType::get(llvm::StructType::create( Ctx, "opencl.event_t"), 0); @@ -76,3 +76,12 @@ return PipeTy; } + +llvm::PointerType *CGOpenCLRuntime::getSamplerType() { + if (!SamplerTy) + SamplerTy = llvm::PointerType::get(llvm::StructType::create( + CGM.getLLVMContext(), "__sampler"), + CGM.getContext().getTargetAddressSpace( + LangAS::opencl_constant)); + return SamplerTy; +} Index: lib/CodeGen/CodeGenModule.h =================================================================== --- lib/CodeGen/CodeGenModule.h +++ lib/CodeGen/CodeGenModule.h @@ -1145,6 +1145,9 @@ llvm::SanitizerStatReport &getSanStats(); + llvm::Value * + createOpenCLIntToSamplerConversion(const Expr *E, CodeGenFunction &CGF); + private: llvm::Constant * GetOrCreateLLVMFunction(StringRef MangledName, llvm::Type *Ty, GlobalDecl D, Index: lib/CodeGen/CodeGenModule.cpp =================================================================== --- lib/CodeGen/CodeGenModule.cpp +++ lib/CodeGen/CodeGenModule.cpp @@ -2382,6 +2382,11 @@ const VarDecl *InitDecl; const Expr *InitExpr = D->getAnyInitializer(InitDecl); + // OpenCL global variables of sampler type are translated to function calls, + // therefore no need to be translated. + if (getLangOpts().OpenCL && ASTTy->isSamplerT()) + return; + // CUDA E.2.4.1 "__shared__ variables cannot have an initialization // as part of their declaration." Sema has already checked for // error cases, so we just need to set Init to UndefValue. @@ -4288,3 +4293,31 @@ return *SanStats; } +llvm::Value * +CodeGenModule::createOpenCLIntToSamplerConversion(const Expr *E, + CodeGenFunction &CGF) { + llvm::Constant *C = EmitConstantExpr(E, E->getType(), &CGF); + + const llvm::ConstantInt *CI = cast(C); + const uint64_t SamplerValue = CI->getValue().getZExtValue(); + // 32-bit value of sampler's initializer is interpreted as + // bit-field with the following structure: + // |unspecified|Filter|Addressing Mode| Normalized Coords| + // |31 6|5 4|3 1| 0| + // This structure corresponds to enum values of sampler properties defined + // in SPIR spec v1.2 and also opencl-c.h + unsigned AddressingMode = (0x0E & SamplerValue) >> 1; + unsigned FilterMode = (0x30 & SamplerValue) >> 4; + if (FilterMode != 1 && FilterMode != 2) + getDiags().Report(Context.getFullLoc(E->getLocStart()), + diag::warn_sampler_initializer_invalid_bits) << "Filter Mode"; + if (AddressingMode > 4) + getDiags().Report(Context.getFullLoc(E->getLocStart()), + diag::warn_sampler_initializer_invalid_bits) << "Addressing Mode"; + + auto SamplerT = getOpenCLRuntime().getSamplerType(); + auto FTy = llvm::FunctionType::get(SamplerT, {C->getType()}, false); + return CGF.Builder.CreateCall(CreateRuntimeFunction(FTy, + "__translate_sampler_initializer"), + {C}); +} Index: lib/Edit/RewriteObjCFoundationAPI.cpp =================================================================== --- lib/Edit/RewriteObjCFoundationAPI.cpp +++ lib/Edit/RewriteObjCFoundationAPI.cpp @@ -1076,6 +1076,7 @@ case CK_CopyAndAutoreleaseBlockObject: case CK_BuiltinFnToFnPtr: case CK_ZeroToOCLEvent: + case CK_IntToOCLSampler: return false; case CK_BooleanToSignedIntegral: Index: lib/Sema/SemaExpr.cpp =================================================================== --- lib/Sema/SemaExpr.cpp +++ lib/Sema/SemaExpr.cpp @@ -7582,6 +7582,11 @@ } } + if (LHSType->isSamplerT() && RHSType->isIntegerType()) { + Kind = CK_IntToOCLSampler; + return Compatible; + } + return Incompatible; } Index: lib/Sema/SemaInit.cpp =================================================================== --- lib/Sema/SemaInit.cpp +++ lib/Sema/SemaInit.cpp @@ -4886,7 +4886,8 @@ QualType DestType, Expr *Initializer) { if (!S.getLangOpts().OpenCL || !DestType->isSamplerT() || - !Initializer->isIntegerConstantExpr(S.getASTContext())) + (!Initializer->isIntegerConstantExpr(S.Context) && + !Initializer->getType()->isSamplerT())) return false; Sequence.AddOCLSamplerInitStep(DestType); @@ -6904,19 +6905,32 @@ } case SK_OCLSamplerInit: { - assert(Step->Type->isSamplerT() && - "Sampler initialization on non-sampler type."); - - QualType SourceType = CurInit.get()->getType(); - + Expr *Init = CurInit.get(); + QualType SourceType = Init->getType(); + // For copy initialization, get the integer literal. if (Entity.isParameterKind()) { - if (!SourceType->isSamplerT()) + if (!SourceType->isSamplerT()) { S.Diag(Kind.getLocation(), diag::err_sampler_argument_required) << SourceType; - } else if (Entity.getKind() != InitializedEntity::EK_Variable) { - llvm_unreachable("Invalid EntityKind!"); + } else if (const DeclRefExpr *DRE = dyn_cast(Init)) { + Init = const_cast(cast(DRE->getDecl())->getInit()); + if (!Init) + break; + Init = cast(Init)->getSubExpr(); + SourceType = Init->getType(); + } } + if (!Init->isConstantInitializer(S.Context, false)) + S.Diag(Kind.getLocation(), + diag::err_sampler_initializer_not_constant); + if (!SourceType->isIntegerType() || + 32 != S.Context.getIntWidth(SourceType)) + S.Diag(Kind.getLocation(), diag::err_sampler_initializer_not_integer) + << SourceType; + + CurInit = S.ImpCastExprToType(Init, S.Context.OCLSamplerTy, + CK_IntToOCLSampler); break; } case SK_OCLZeroEvent: { Index: lib/StaticAnalyzer/Core/ExprEngineC.cpp =================================================================== --- lib/StaticAnalyzer/Core/ExprEngineC.cpp +++ lib/StaticAnalyzer/Core/ExprEngineC.cpp @@ -341,6 +341,7 @@ case CK_AnyPointerToBlockPointerCast: case CK_ObjCObjectLValueCast: case CK_ZeroToOCLEvent: + case CK_IntToOCLSampler: case CK_LValueBitCast: { // Delegate to SValBuilder to process. SVal V = state->getSVal(Ex, LCtx); Index: test/CodeGenOpenCL/opencl_types.cl =================================================================== --- test/CodeGenOpenCL/opencl_types.cl +++ test/CodeGenOpenCL/opencl_types.cl @@ -1,39 +1,43 @@ -// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s +// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -o - -O0 | FileCheck %s -constant sampler_t glb_smp = 7; -// CHECK: constant i32 7 +#define CLK_ADDRESS_CLAMP_TO_EDGE 2 +#define CLK_NORMALIZED_COORDS_TRUE 1 +#define CLK_FILTER_NEAREST 0x10 +#define CLK_FILTER_LINEAR 0x20 + +constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_NEAREST; void fnc1(image1d_t img) {} -// CHECK: @fnc1(%opencl.image1d_ro_t* +// CHECK: @fnc1(%opencl.image1d_ro_t addrspace(1)* void fnc1arr(image1d_array_t img) {} -// CHECK: @fnc1arr(%opencl.image1d_array_ro_t* +// CHECK: @fnc1arr(%opencl.image1d_array_ro_t addrspace(1)* void fnc1buff(image1d_buffer_t img) {} -// CHECK: @fnc1buff(%opencl.image1d_buffer_ro_t* +// CHECK: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(1)* void fnc2(image2d_t img) {} -// CHECK: @fnc2(%opencl.image2d_ro_t* +// CHECK: @fnc2(%opencl.image2d_ro_t addrspace(1)* void fnc2arr(image2d_array_t img) {} -// CHECK: @fnc2arr(%opencl.image2d_array_ro_t* +// CHECK: @fnc2arr(%opencl.image2d_array_ro_t addrspace(1)* void fnc3(image3d_t img) {} -// CHECK: @fnc3(%opencl.image3d_ro_t* +// CHECK: @fnc3(%opencl.image3d_ro_t addrspace(1)* void fnc4smp(sampler_t s) {} -// CHECK-LABEL: define {{.*}}void @fnc4smp(i32 +// CHECK-LABEL: define {{.*}}void @fnc4smp(%__sampler addrspace(2)* kernel void foo(image1d_t img) { - sampler_t smp = 5; - // CHECK: alloca i32 + sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR; + // CHECK: alloca %__sampler addrspace(2)* event_t evt; // CHECK: alloca %opencl.event_t* - // CHECK: store i32 5, + // CHECK: store %__sampler addrspace(2)* fnc4smp(smp); - // CHECK: call {{.*}}void @fnc4smp(i32 + // CHECK: call {{.*}}void @fnc4smp(%__sampler addrspace(2)* fnc4smp(glb_smp); - // CHECK: call {{.*}}void @fnc4smp(i32 + // CHECK: call {{.*}}void @fnc4smp(%__sampler addrspace(2)* } void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {} Index: test/CodeGenOpenCL/sampler.cl =================================================================== --- /dev/null +++ test/CodeGenOpenCL/sampler.cl @@ -0,0 +1,47 @@ +// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 -Wspir-compat -verify -DCHECK_SAMPLER_VALUE | FileCheck %s +// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 -verify | FileCheck %s + +#ifndef CHECK_SAMPLER_VALUE +// expected-no-diagnostics +#endif + +#define CLK_ADDRESS_CLAMP_TO_EDGE 2 +#define CLK_NORMALIZED_COORDS_TRUE 1 +#define CLK_FILTER_NEAREST 0x10 +#define CLK_FILTER_LINEAR 0x20 + +constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR; + +constant sampler_t glb_smp2 = 0; +#ifdef CHECK_SAMPLER_VALUE +// expected-warning@-2{{sampler initializer has invalid Filter Mode bits}} +#endif + +constant sampler_t glb_smp3 = 0x1f; +#ifdef CHECK_SAMPLER_VALUE +// expected-warning@-2{{sampler initializer has invalid Addressing Mode bits}} +#endif + +// CHECK: %__sampler = type opaque + +void fnc4smp(sampler_t s) {} +// CHECK: define spir_func void @fnc4smp(%__sampler addrspace(2)* % + +kernel void foo() { + sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST; + // CHECK-LABEL: define spir_kernel void @foo() + // CHECK: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca %__sampler addrspace(2)* + // CHECK: [[SAMP:%[0-9]+]] = call %__sampler addrspace(2)* @__translate_sampler_initializer(i32 19) + // CHECK: store %__sampler addrspace(2)* [[SAMP]], %__sampler addrspace(2)** [[smp_ptr]] + + fnc4smp(smp); + // CHECK: [[SAMP:%[0-9]+]] = call %__sampler addrspace(2)* @__translate_sampler_initializer(i32 19) + // CHECK: call spir_func void @fnc4smp(%__sampler addrspace(2)* [[SAMP]]) + + fnc4smp(glb_smp); + // CHECK: [[SAMP:%[0-9]+]] = call %__sampler addrspace(2)* @__translate_sampler_initializer(i32 35) + // CHECK: call spir_func void @fnc4smp(%__sampler addrspace(2)* [[SAMP]]) + + fnc4smp(glb_smp2); + fnc4smp(glb_smp3); +} Index: test/SemaOpenCL/sampler_t.cl =================================================================== --- test/SemaOpenCL/sampler_t.cl +++ test/SemaOpenCL/sampler_t.cl @@ -13,6 +13,7 @@ const sampler_t const_smp = 7; foo(glb_smp); foo(const_smp); + foo(argsmp); foo(5); // expected-error {{sampler_t variable required - got 'int'}} sampler_t sa[] = {argsmp, const_smp}; // expected-error {{array of 'sampler_t' type is invalid in OpenCL}} }