yaxunl updated the summary for this revision.
yaxunl updated this revision to Diff 63286.
yaxunl added a comment.
Drop the sampler initializer structure as Anastasia suggested. Now
`__translate_sampler_initializer` accepts an integer parameter.
http://reviews.llvm.org/D21567
Files:
include/clang/AST/OperationKinds.def
include/clang/Basic/DiagnosticGroups.td
include/clang/Basic/DiagnosticSemaKinds.td
lib/AST/ASTContext.cpp
lib/AST/Expr.cpp
lib/AST/ExprConstant.cpp
lib/CodeGen/CGDebugInfo.cpp
lib/CodeGen/CGDebugInfo.h
lib/CodeGen/CGExpr.cpp
lib/CodeGen/CGExprAgg.cpp
lib/CodeGen/CGExprComplex.cpp
lib/CodeGen/CGExprConstant.cpp
lib/CodeGen/CGExprScalar.cpp
lib/CodeGen/CGOpenCLRuntime.cpp
lib/CodeGen/CGOpenCLRuntime.h
lib/CodeGen/CodeGenModule.cpp
lib/CodeGen/CodeGenModule.h
lib/Edit/RewriteObjCFoundationAPI.cpp
lib/Sema/SemaExpr.cpp
lib/Sema/SemaInit.cpp
lib/StaticAnalyzer/Core/ExprEngineC.cpp
test/CodeGenOpenCL/opencl_types.cl
test/CodeGenOpenCL/sampler.cl
test/SemaOpenCL/sampler_t.cl
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}}
}
Index: test/CodeGenOpenCL/sampler.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/sampler.cl
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+
+#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;
+
+// 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]])
+}
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: 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: 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,33 @@
}
case SK_OCLSamplerInit: {
- assert(Step->Type->isSamplerT() &&
- "Sampler initialization on non-sampler type.");
-
- QualType SourceType = CurInit.get()->getType();
-
+ Expr *Init = CurInit.get();
+ Init->dump();
+ 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<DeclRefExpr>(Init)) {
+ Init = const_cast<Expr*>(cast<VarDecl>(DRE->getDecl())->getInit());
+ if (!Init)
+ break;
+ Init = cast<ImplicitCastExpr>(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/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/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/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<llvm::ConstantInt>(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/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/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/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/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/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/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/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/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/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/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: 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<SpirCompat>;
+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: 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/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.
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
