[OpenCL] Generate struct type for sampler_t and function call for the initializer

[yaxunl]

Summary

Currently Clang use int32 to represent sampler_t, which have been a source of issue for some backends, because in some backends sampler_t cannot be represented by int32. They have to depend on kernel argument metadata and use IPA to find the sampler arguments and global variables and transform them to target specific sampler type.

This patch uses opaque pointer type opencl.sampler_t* for sampler_t. For each use of program-scope sampler variable, it generates a function call of __translate_sampler_initializer. For each initialization of function-scope sampler variable, it generates a function call of __translate_sampler_initializer. For example,

constant sampler_t s1 = X1; // X1 is an integer literal
 
void f() {
  f1(s1);
  f2(s1);
  const sampler_t s2 = X2; // X2 is an integer literal
  f3(s2);
  f4(s2);
}

=> Llvm bitcode equivalent to (assuming opencl.sampler_t is the opaque type to represent sampler_t):

// opaque struct type for sampler
struct opencl.sampler_t;

constant opencl.sampler_t *__attribute__((always_inline)) __translate_sampler_initializer(unsigned); // a builtin function for translating sampler literal
 
void f() {
  f1(__translate_sampler_initializer(X1));
  f2(__translate_sampler_initializer(X1));
  constant __opencl_sampler_t *s2 = __translate_sampler_initializer(X2);
  f3(s2);
  f4(s2);
}

Each builtin library can implement its own __translate_sampler_initializer(). Since the real sampler type tends to be architecture dependent, allowing it to be initialized by a library function simplifies backend design. A typical implementation of __translate_sampler_initializer could be a table lookup of real sampler literal values. Since its argument is always a literal, the returned pointer is known at compile time and easily optimized to finally become some literal values directly put into image read instructions.

The advantage of this representation is:

1. Robust - can be optimized without losing information
2. Easy to implement – can be implemented by library instead of pass

The implementation of the design is to introduce a opaque sampler type and translates each sampler variable to sampler type in AST, and introduces cast from integer to sampler for each variable initializer. For copy initialization with function-scope variable, create integer to sampler cast. In codegen, sampler type is translated to opaque opencl.sampler_t* type. The cast from integer to sampler is translated to function call __translate_sampler_initializer. The global variable of sampler type is not translated since they are no longer referenced.

This patch is partially based on Alexey Sotkin's work in Khronos Clang (https://github.com/KhronosGroup/SPIR/commit/3d4eec61623502fc306e8c67c9868be2b136e42b).

[yaxunl]

Add test deleted by accident.

[Anastasia]

include/clang/AST/BuiltinTypes.def

164 ↗

(On Diff #61443)

I can't get why is this necessary to represent initializer as a special Clang type? Could we avoid this additional complexity?

include/clang/AST/OperationKinds.def

328

Could we just have only int->sampler conversion? Without having an extra type for initializer?

include/clang/Basic/DiagnosticGroups.td

876

I don't understand the description really? Why not compatible?

include/clang/Driver/CC1Options.td

690 ↗

(On Diff #61443)

Any reason to have this flag and support different sampler representations in Clang?

lib/CodeGen/CGOpenCLRuntime.cpp

52

Could you please keep coherency in type naming i.e. add "opencl." prefix.

[yaxunl]

include/clang/AST/BuiltinTypes.def

164 ↗

(On Diff #61443)

If we don't change the type in AST, we need to translate the sampler type in AST to different types based on whether the entry is a variable or a function argument. This needs some ugly hacking of the codegen.

include/clang/AST/OperationKinds.def

328

If we don't insert the sampler initializer to sampler cast in AST, when we want to insert the function call __translate_sampler_initializer in codegen, we have to hacking the generic translation of function call instruction and other instructions which may have a sampler variable as operand to insert the function call, which would be ugly.

Or we could do a post-processing of the llvm module at the end of codegen to insert the function calls, if this approach is acceptable.

include/clang/Basic/DiagnosticGroups.td

876

e.g. the sampler constant may take some literal value which is not compatible with SPIR/SPIR-V spec. A warning will be emitted. Such warning msg belongs to this category.

include/clang/Driver/CC1Options.td

690 ↗

(On Diff #61443)

Give user some time for transition.

lib/CodeGen/CGOpenCLRuntime.cpp

52

We need to be able to implement function

__sampler* __translate_sampler_initializer(__sampler_initializer*);

in library, where __sampler and __sampler_initializer are concrete struct types defined in the library source code. Therefore we cannot have dot in the struct type name.

[Anastasia]

include/clang/AST/BuiltinTypes.def

164 ↗

(On Diff #61443)

I don't understand why this is needed. Surely, sampler type is just a sampler type independently from the context it's being used. Could you elaborate or give me an example please.

include/clang/AST/OperationKinds.def

328

I don't understand why this is needed? The initialization function only replaces the initializer of a sampler type. So this call only has to be emitted on the conversion int->sampler.

include/clang/Basic/DiagnosticGroups.td

876

Could we avoid having SPIR incompatible code instead?

include/clang/Driver/CC1Options.td

690 ↗

(On Diff #61443)

It doesn't look such a big and risky change to add special flag and maintain two different representations. I would rather go for cleanup of old code here.

lib/CodeGen/CGOpenCLRuntime.cpp

52

Would it work if this function uses OpenCL original types and get compiled with Clang too?

[yaxunl]

include/clang/AST/BuiltinTypes.def

164 ↗

(On Diff #61443)

For example,

constant sampler_t a = 0;
kernel void f(sampler_t b) {
   g(a);
   g(b);
 }

Let's say we don't introduce a OCLSamplerInitTy in AST. We only introduce OCLSamplerTy to represent sampler_t in AST. Then both variables a and b have OCLSamplerTy type in AST.

In codegen, type translation is done by CodeGenTypes::ConvertType, which only accepts a QualType argument, which means we can only translate an AST type to an LLVM type based on the AST type itself. Therefore we can only translate OCLSamplerTy to one LLVM type. Let's say we translate OCLSamplerTy to __sampler* type. Then variable a will be translated to __sampler* type. However, we want it to be translated to __sampler_initializer* type. To do that, we have to customize CodeGenModule::EmitGlobalVarDefinition to translate global variables with OCLSamplerTy type specially, which will be ugly.

Basically we cannot avoid the complexity of translating the type for variable a. The complexity is only moved from AST to Codegen. If we introduce OCLSamplerInitTy in AST, we don't need to translate the type of variable a specially. It seems to be a more elegant solution to me.

include/clang/AST/OperationKinds.def

328

Let's consider this example:

constant sampler_t a = 0;
kernel void f() {
   g(a);
 }

In codegen, we want to insert a call of __transform_sampler_initializer for the reference of variable a, not the definition of variable a. If we introduce a cast SamplerInitializerToSamplerCast, we can represent that in AST equivalent to:

__sampler_initializer *a = IntToSamplerInitializerCast(0);
kernel void f() {
   g(SamplerInitializerToSamplerCast(a));
}

We can do a straight forward translation by translating SamplerInitializerToSamplerCast to function call of __transform_sampler_initializer.

If we do not have SamplerInitializerToSamplerCast, the AST will be equivalent to:

__sampler *a = IntToSamplerCast(0);
kernel void f() {
   g(a);
}

Can we translate IntToSamplerCast to __transform_sampler_initializer? No. Because translation of IntToSamplerCast is done at translating global variable a, which happens before translating instructions belonging to a function, whereas we need to insert a function call of __transform_sampler_initializer when translating function call of g. Since we do not have a landmark for inserting function call of __transform_sampler_initializer, we have to customize CodeGenFunction::EmitCall to handle argument being a global variable of type __sampler specially, which will be ugly.

However we do have a simple solution for this if we do a post-processing of the generated LLVM IR in CodeGenModule::Release. Basically we iterate all uses of global variable a and replace them with a function call of __transform_sampler_initializer. By doing this we can use a simple AST without introducing the sampler initializer type and the cast from sampler initializer to sampler. Since CodeGenModule::Release already contains quite a few post-prcessing's for LLVM IR, I think it is reasonable to do some OpenCL specific post-processing there. We can introduce a new member CGOpenCLRuntime::postprocessIR for this purpose.

include/clang/Basic/DiagnosticGroups.td

876

It is the user's OpenCL source code which may contain SPIR incompatible integer literal value for sampler, e.g. a user may write

sampler_t s = 0;

This will result in a sampler value not compatible with SPIR. We cannot prevent user from doing that, but we can emit a warning for that.

include/clang/Driver/CC1Options.td

690 ↗

(On Diff #61443)

I can remove the support of old sampler representation.

lib/CodeGen/CGOpenCLRuntime.cpp

52

Probably not. The OpenCL type sampler_t is translated to a pointer to an opaque struct. However we need to return a pointer to a concrete struct type since we need to fill the struct in __transform_sampler_initializer.

[yaxunl]

Removed the old representation of sampler type by i32.

[Anastasia]

include/clang/AST/BuiltinTypes.def

164 ↗

(On Diff #62966)

However, we want it to be translated to __sampler_initializer* type.

This is precisely what I am not able to understand. Why do we want it to be __sampler_initializer* type? I don't think we want that. If it was a sampler type originally in OpenCL, it should be kept a sampler type in LLVM too.

include/clang/AST/OperationKinds.def

328

In codegen, we want to insert a call of __transform_sampler_initializer for the reference of variable a, not the definition of variable a.

I think the original discussion was to insert the call on the declaration of sampler variable as an initializer. This would allow us to use variable itself "as is" everywhere else.

So if let's say we compile:

constant sampler_t a = 0;
kernel void f() {
   g(a);
}

the code we are generating would be equivalent to:

kernel void f() {
  __sampler* a = __transform_sampler_initializer(0);
  g(a);
}

Why are we not using this approach?

I think we discussed later that we could generate a struct of initializers instead of just int (0 -> struct sampler_initializer). Here is the description: http://lists.llvm.org/pipermail/cfe-dev/2016-June/049389.html
But general principle still was as described above...

include/clang/Basic/DiagnosticGroups.td

876

I see. This is just because OpenCL doesn't allocate actual constants for different allowed sampler modes whereas SPIR does.

lib/CodeGen/CGOpenCLRuntime.cpp

52

If you do conversion struct ptr <-> opaque ptr in this function, it should be fine? I would imagine images and other OpenCL types have the same requirement.

[yaxunl]

+ Brian

Hi Anastasia,

The advantage for translating sampler variable to a global variable with __sampler_initializer type is that it is consistent with OpenCL C++ and SPIRV, so it is easier to translate the IR to SPIRV.

About the type name of sampler_t in LLVM, using a name without . allows library functions to use the concrete sampler types directly without casting. Casting not only affects the readability of the library code, but also causes extra efforts in optimizing these codes.

Sam

[yaxunl]

Drop the sampler initializer structure as Anastasia suggested. Now __translate_sampler_initializer accepts an integer parameter.

[Anastasia]

lib/CodeGen/CodeGenModule.cpp

4309

Seems like this code is specific to SPIR and not OpenCL?

lib/Sema/SemaInit.cpp

6907

Is the assert no longer needed?

6911

Debug output?

test/CodeGenOpenCL/sampler.cl

24

Why does the initialization happen second time here?

27

Could we check definition of global sampler variable in IR too?

[Anastasia]

include/clang/Basic/DiagnosticSemaKinds.td

7868

start in lower case as all other messages!

lib/CodeGen/CGExprConstant.cpp

694

Does this only apply to global variable samplers?

[yaxunl]

lib/CodeGen/CGExprConstant.cpp

694

Yes. The translation of local sampler is like

sampler_t p = X;

->

__sampler** p = alloca __sampler*;
store __translate_sampler_initializer(X), p

lib/CodeGen/CodeGenModule.cpp

4309

This warning msg is emitted for SPIR incompatible sampler value. It belongs to a category spir-compat. How about we turn off this category of warning by default and only enable it for SPIR target by default?

lib/Sema/SemaInit.cpp

6907

This assert is unlikely to happen. not so useful.

test/CodeGenOpenCL/sampler.cl

24

We cannot call function in the initializer of an LLVM global variable, so each reference of the global variable is replaced by a call of __transform_sampler_initializer and the original global variable is eliminated.

27

The original global variable is eliminated and its references are replaced by call of __transform_sampler_initializer.

[yaxunl]

Revised by Anastia's comments. Disabled warning for sampler enum value by default.

[yaxunl]

Added test for invalid sampler values.

[Anastasia]

lib/CodeGen/CGOpenCLRuntime.cpp

83

I am wondering if we could at least name it __opencl_sampler_t. I kind of don't like that we have inconsistency with the rest here.

Btw, the types in Debug are all named opencl_<type>. Perhaps we could correct them consistently across?

lib/CodeGen/CodeGenModule.cpp

4309

Yes sounds good!

lib/Sema/SemaInit.cpp

6907

But I am guessing that holds for most of the assertions. :)

6917–6929

Should you also check for type of Expr to be sampler?

Btw is this covered in testing?

test/CodeGenOpenCL/sampler.cl

2

I think we can give this warning for spir target or -Wspir-compat.

18

This should go into SemaOpenCL tests!

28

Could we add CHECK-NOT respectively?

38

Can we add another call with smp and check that __translate_sampler_initializeris is not called again?

[yaxunl]

lib/CodeGen/CGOpenCLRuntime.cpp

83

I will rename it as __opencl_sampler_t.

I would suggest to rename all opencl types as __opencl_x instead of opencl.x (in another patch) to avoid unnecessary casts in libraries.

lib/Sema/SemaInit.cpp

6907

I can add it back.

6917–6929

The type of Init is SourceType. We already checked it in line 6912.

Here we want to handle the case where a sampler type variable is passed as function argument. Since DRE->getDecl() has the same type as DRE, so we don't need to check that.

Since we don't allow operators on sampler, the only sampler expressions we can get are

1. sampler type global variable
2. sampler type function-scope variable
3. sampler type function parameter

I covered the 3 cases in the tests. Also covered passing non-sampler type expression as sampler argument to a function. However I noticed I did not cover the cases that sampler type variable without initializer. I will add tests for that.

test/CodeGenOpenCL/sampler.cl

2

Sorry I did not turn on the warning for spir-compat by default for SPIR target. As I checked the part of code about warnings, I did not find a precedence to turn on warning for a specific target? So here I still need -Wspir-compat to see the warnings for SPIR target. Should I turn on the warning for spir-compat by default for SPIR target.

18

This warning is emitted at codegen stage instead of parsing since diagnosing this requires evaluated value of the initializer expression, which is readily available at codegen, but need some extra effort to obtain at parsing.

One side effect is that we will only see the warning if the test does not contain other sema errors because otherwise we will not reach codegen stage. That's why I test it in a codegen test since the sema tests contain other errors.

28

Will do.

38

Will do. Currently it generates call for each reference of smp, but I will fix that.

[yaxunl]

Revised by Anastasia's comments. This fixes multiple function call of __translate_sampler_initializer for function-scope sampler variables.

[yaxunl]

Turn on -Wspir-compat for spir target by default. Diagnose invalid sampler value at Sema instead of Codegen.

[yaxunl]

test/CodeGenOpenCL/sampler.cl

2

Fixed and moved to sema test.

18

Moved to sema tests.

[Anastasia]

lib/CodeGen/CodeGenModule.cpp

2399

Could we lift this check right to the beginning of the function?

lib/Sema/SemaInit.cpp

6917–6929

It would be nice to put some comments here summarizing your description!

Btw, I don't think it's covered in testing yet, isn't it?

6918

I don't get why this level of indirection is added? Should the variable initialization be handled elsewhere?

6922

CurInit.get() can be changed to Init

test/SemaOpenCL/sampler_t.cl

55

Why was this change required?

[yaxunl]

Revised by Anastasia's comments.

lib/Sema/SemaInit.cpp

6917–6929

Added detailed comments.

I only missed passing sampler parameter as function call argument. Now added.

6918

For global variables, since we cannot initialize them with a function call __translate_sampler_initializer, we have to take their initializer and pass them to a function call.

test/SemaOpenCL/sampler_t.cl

55

Since i added a run with spir triple, there is an extra error msg emitted:

function with no prototype cannot use the spir_function calling convention

adding void argument suppresses this error msg.

[Anastasia]

lib/Sema/SemaInit.cpp

6917–6959

Is this being tested anywhere:

sampler_t initialization requires 32-bit integer

6930

What if global variable sampler is initialized with another sampler variable:

sampler_t s1 = ...;
sampler_t s2 = s1;
...
foo(s2);

Btw, I am wondering whether this code is needed at all, because I am guessing variable initialization will be handled separately anyways, irrespective to whether it's being used in a call or not...

[yaxunl]

lib/Sema/SemaInit.cpp

6930

clang currently does not allow assigning a sampler global variable with another sampler global variable. An error will be emitted:

initializer element is not a compile-time constant

We could allow this, but I think allowing this is not a very useful feature, so I would recommend continue not allowing it.

If we do not replace references of function-scope sampler variable with its initializer here, we will end up with a global variable initialized with an integer and we will not be able to do codegen for it as in LLVM we cannot initialize a global variable with a function call.

[yaxunl]

Added more diagnostic tests.

[yaxunl]

Update the test to accommodate changes in branch tip.

[Anastasia]

lib/Sema/SemaInit.cpp

6930

Could you please take a look at the new diagnostics you are adding in DiagnosticSemaKinds.td. I don't think they are all covered by your tests.

I understand that you have to insert the cast even though it looks a bit hacky, so we might need to put more comments here to explain why this is needed. But certainly diagnostics code seems to be doing extra work. For the following code:

constant sampler_t glb_smp = 1L;
void foo(sampler_t);
void bar(){
  foo(glb_smp);
}

I get diagnostics reported twice:

sampler.cl:1:20: error: sampler_t initialization requires 32-bit integer, not 'long'
constant sampler_t glb_smp = 1L;
sampler.cl:6:5: error: sampler_t initialization requires 32-bit integer, not 'long'
foo(glb_smp);

The second diagnostic doesn't seem right at all.

Could you please take a look.

[yaxunl]

Add more sema tests. Eliminate redundant diagnostic messages.

[Anastasia]

Btw, I am missing tests for generated __translate_sampler_initializer which I think we had at some point.

lib/CodeGen/CGOpenCLRuntime.cpp

83

Related to this thread: http://lists.llvm.org/pipermail/cfe-dev/2016-July/050126.html.

I am guessing leaving the old naming scheme should be fine.

lib/Sema/SemaInit.cpp

6907

There are 5 cases it seems :)

6931

beend -> been

6933

Do we accept declaring global sampler_t variable without an initializer?

6940

Could you use else here instead of adding TakeLiteral variable to avoid going to the next compound statement?

6949

Should this generate an error?

test/SemaOpenCL/sampler_t.cl

35

Are these extra calls below adding any extra testing?

53

Does this line add extra testing?

[yaxunl]

lib/CodeGen/CGOpenCLRuntime.cpp

83

I have some more comments about that. I think we should use struct.__opencl_sampler as the name of sampler type in LLVM IR.

lib/Sema/SemaInit.cpp

6933

no. Since in file scope sampler variable is only allowed in constant addr space, it has to be initialized. There is a test for it.

6940

Good catch. Will do.

6949

This has already been diagnosed when parsing the intializer of the variable declaration.

test/SemaOpenCL/sampler_t.cl

35

Add the function calls to test

1. No additional error msgs since the error msgs should be emitted for the variable declarations.
2. no crashes when parsing them

53

Test no crashes when an invalid sampler expression is passed as a sampler argument.

[yaxunl]

Revised by Anastasia's comments. Added missing codegen test.

[yaxunl]

Use 'opencl.sampler_t' as opque type name for sampler in LLVM IR.

[Anastasia]

LGTM! Thanks! Could you please address these small comments before committing!

lib/Sema/SemaInit.cpp

6944

Can you remove this sentence as it duplicates the previous one.

6946

I think you don't need this check any more because this code is inside the else part. Could you please double check before committing?

test/CodeGenOpenCL/sampler.cl

4

5 cases

[yaxunl]

lib/Sema/SemaInit.cpp

6946

I tried removing this line. There will be extra diagnostics emitted for test case:

constant sampler_t glb_smp6 = glb_smp;

We only expects error msg: initializer element is not a compile-time constant. However in addition to that, we also get sampler_t initialization requires 32-bit integer, not '__constant sampler_t'. This is because the break is removed. So I think better not removed the line.

[Anastasia]

lib/Sema/SemaInit.cpp

6946

Sure! Thanks for checking it!

include/clang/AST/OperationKinds.def

 CAST_OPERATION(BuiltinFnToFnPtr)
 
 // Convert a zero value for OpenCL event_t initialization.
 CAST_OPERATION(ZeroToOCLEvent)
 
 // 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.
 // Note that additions to this should also update the StmtVisitor class.
 
 // [C++ 5.5] Pointer-to-member operators.
 BINARY_OPERATION(PtrMemD, ".*")
 BINARY_OPERATION(PtrMemI, "->*")

include/clang/Basic/DiagnosticGroups.td

 // A warning group for things that will change semantics in the future.
 def FutureCompat : DiagGroup<"future-compat">;
 
 def InvalidOrNonExistentDirectory : DiagGroup<"invalid-or-nonexistent-directory">;
 
 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

include/clang/Basic/DiagnosticSemaKinds.td

   "field of illegal %select{type|pointer type}0 %1 declared here">;
 def err_event_t_global_var : Error<
   "the event_t type cannot be used to declare a program scope variable">;
 def err_opencl_type_struct_or_union_field : Error<
   "the %0 type cannot be used to declare a structure or union field">;
 def err_event_t_addr_space_qual : Error<
   "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>, 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<
   "sampler type cannot be used with the __local and __global address space qualifiers">;
 def error_opencl_cast_non_zero_to_event_t : Error<
   "cannot cast non-zero value '%0' to 'event_t'">;
 def err_opencl_global_invalid_addr_space : Error<
   "%select{program scope|static local|extern}0 variable must reside in %1 address space">;

lib/AST/ASTContext.cpp

       Align = Target->getPointerAlign(0); //   == sizeof(void*)
       break;
     case BuiltinType::ObjCId:
     case BuiltinType::ObjCClass:
     case BuiltinType::ObjCSel:
       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:
     case BuiltinType::OCLNDRange:
     case BuiltinType::OCLReserveID:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"

lib/AST/Expr.cpp

   case CK_IntegralComplexToReal:
   case CK_IntegralComplexCast:
   case CK_IntegralComplexToFloatingComplex:
   case CK_ARCProduceObject:
   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:
   case CK_ARCExtendBlockObject:
   case CK_ZeroToOCLEvent:
+  case CK_IntToOCLSampler:
     assert(!getType()->isBooleanType() && "unheralded conversion to bool");
     goto CheckNoBasePath;
 
   case CK_Dependent:
   case CK_LValueToRValue:
   case CK_NoOp:
   case CK_AtomicToNonAtomic:
   case CK_NonAtomicToAtomic:

     const CastExpr *CE = cast<CastExpr>(this);
 
     // Handle misc casts we want to ignore.
     if (CE->getCastKind() == CK_NoOp ||
         CE->getCastKind() == CK_LValueToRValue ||
         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;
   }
   case MaterializeTemporaryExprClass:
     return cast<MaterializeTemporaryExpr>(this)->GetTemporaryExpr()
       ->isConstantInitializer(Ctx, false, Culprit);
 

lib/AST/ExprConstant.cpp

   case CK_FloatingComplexToIntegralComplex:
   case CK_IntegralRealToComplex:
   case CK_IntegralComplexCast:
   case CK_IntegralComplexToFloatingComplex:
   case CK_BuiltinFnToFnPtr:
   case CK_ZeroToOCLEvent:
   case CK_NonAtomicToAtomic:
   case CK_AddressSpaceConversion:
+  case CK_IntToOCLSampler:
     llvm_unreachable("invalid cast kind for integral value");
 
   case CK_BitCast:
   case CK_Dependent:
   case CK_LValueBitCast:
   case CK_ARCProduceObject:
   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:

   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:
   case CK_ARCExtendBlockObject:
   case CK_CopyAndAutoreleaseBlockObject:
   case CK_BuiltinFnToFnPtr:
   case CK_ZeroToOCLEvent:
   case CK_NonAtomicToAtomic:
   case CK_AddressSpaceConversion:
+  case CK_IntToOCLSampler:
     llvm_unreachable("invalid cast kind for complex value");
 
   case CK_LValueToRValue:
   case CK_AtomicToNonAtomic:
   case CK_NoOp:
     return ExprEvaluatorBaseTy::VisitCastExpr(E);
 
   case CK_Dependent:

lib/CodeGen/CGDebugInfo.h

class CGDebugInfo {
   SourceLocation CurLoc;
   llvm::DIType *VTablePtrType = nullptr;
   llvm::DIType *ClassTy = nullptr;
   llvm::DICompositeType *ObjTy = nullptr;
   llvm::DIType *SelTy = nullptr;
 #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;
   llvm::DIType *OCLNDRangeDITy = nullptr;
   llvm::DIType *OCLReserveIDDITy = nullptr;
 
   /// Cache of previously constructed Types.
   llvm::DenseMap<const void *, llvm::TrackingMDRef> TypeCache;

lib/CodeGen/CGDebugInfo.cpp

#include "clang/AST/BuiltinTypes.def"
   }
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     return getOrCreateStructPtrType("opencl_" #ImgType "_" #Suffix "_t", \
                                     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:
     return getOrCreateStructPtrType("opencl_clk_event_t", OCLClkEventDITy);
   case BuiltinType::OCLQueue:
     return getOrCreateStructPtrType("opencl_queue_t", OCLQueueDITy);
   case BuiltinType::OCLNDRange:
     return getOrCreateStructPtrType("opencl_ndrange_t", OCLNDRangeDITy);

lib/CodeGen/CGExpr.cpp

   case CK_ReinterpretMemberPointer:
   case CK_AnyPointerToBlockPointerCast:
   case CK_ARCProduceObject:
   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:
   case CK_ARCExtendBlockObject:
   case CK_CopyAndAutoreleaseBlockObject:
   case CK_AddressSpaceConversion:
+  case CK_IntToOCLSampler:
     return EmitUnsupportedLValue(E, "unexpected cast lvalue");
 
   case CK_Dependent:
     llvm_unreachable("dependent cast kind in IR gen!");
 
   case CK_BuiltinFnToFnPtr:
     llvm_unreachable("builtin functions are handled elsewhere");
 

lib/CodeGen/CGExprAgg.cpp

void AggExprEmitter::VisitCastExpr(CastExpr *E) {
   case CK_ARCProduceObject:
   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:
   case CK_ARCExtendBlockObject:
   case CK_CopyAndAutoreleaseBlockObject:
   case CK_BuiltinFnToFnPtr:
   case CK_ZeroToOCLEvent:
   case CK_AddressSpaceConversion:
+  case CK_IntToOCLSampler:
     llvm_unreachable("cast kind invalid for aggregate types");
   }
 }
 
 void AggExprEmitter::VisitCallExpr(const CallExpr *E) {
   if (E->getCallReturnType(CGF.getContext())->isReferenceType()) {
     EmitAggLoadOfLValue(E);
     return;

lib/CodeGen/CGExprComplex.cpp

ComplexPairTy ComplexExprEmitter::EmitCast(CastKind CK, Expr *Op,
   case CK_ARCProduceObject:
   case CK_ARCConsumeObject:
   case CK_ARCReclaimReturnedObject:
   case CK_ARCExtendBlockObject:
   case CK_CopyAndAutoreleaseBlockObject:
   case CK_BuiltinFnToFnPtr:
   case CK_ZeroToOCLEvent:
   case CK_AddressSpaceConversion:
+  case CK_IntToOCLSampler:
     llvm_unreachable("invalid cast kind for complex value");
 
   case CK_FloatingRealToComplex:
   case CK_IntegralRealToComplex:
     return EmitScalarToComplexCast(CGF.EmitScalarExpr(Op), Op->getType(),
                                    DestTy, Op->getExprLoc());
 
   case CK_FloatingComplexCast:

lib/CodeGen/CGExprConstant.cpp

llvm::Constant *VisitCastExpr(CastExpr* E) {
 
     case CK_LValueToRValue:
     case CK_AtomicToNonAtomic:
     case CK_NonAtomicToAtomic:
     case CK_NoOp:
     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:
       llvm_unreachable("builtin functions are handled elsewhere");
 
     case CK_ReinterpretMemberPointer:
     case CK_DerivedToBaseMemberPointer:
     case CK_BaseToDerivedMemberPointer:

lib/CodeGen/CGExprScalar.cpp

                                          CE->getExprLoc());
   }
 
   case CK_ZeroToOCLEvent: {
     assert(DestTy->isEventT() && "CK_ZeroToOCLEvent cast on non-event type");
     return llvm::Constant::getNullValue(ConvertType(DestTy));
   }
 
-  }
+  case CK_IntToOCLSampler:
+    return CGF.CGM.createOpenCLIntToSamplerConversion(E, CGF);
+
+  } // end of switch
 
   llvm_unreachable("unknown scalar cast");
 }
 
 Value *ScalarExprEmitter::VisitStmtExpr(const StmtExpr *E) {
   CodeGenFunction::StmtExprEvaluation eval(CGF);
   Address RetAlloca = CGF.EmitCompoundStmt(*E->getSubStmt(),
                                            !E->getType()->isVoidType());

lib/CodeGen/CGOpenCLRuntime.h

 
 class CodeGenFunction;
 class CodeGenModule;
 
 class CGOpenCLRuntime {
 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
   /// an entry to CGF's LocalDeclMap for D.  The base class does this using
   /// CodeGenFunction::EmitStaticVarDecl to emit an internal global for D.
   virtual void EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF,
                                          const VarDecl &D);
 
   virtual llvm::Type *convertOpenCLSpecificType(const Type *T);
 
   virtual llvm::Type *getPipeType();
+
+  llvm::PointerType *getSamplerType();
 };
 
 }
 }
 
 #endif

lib/CodeGen/CGOpenCLRuntime.cpp

default:
     return nullptr;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     return llvm::PointerType::get( \
         llvm::StructType::create(Ctx, "opencl." #ImgType "_" #Suffix "_t"), \
         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);
   case BuiltinType::OCLClkEvent:
     return llvm::PointerType::get(
         llvm::StructType::create(Ctx, "opencl.clk_event_t"), 0);
   case BuiltinType::OCLQueue:
     return llvm::PointerType::get(
         llvm::StructType::create(Ctx, "opencl.queue_t"), 0);
   case BuiltinType::OCLNDRange:
if (!PipeTy){
     uint32_t PipeAddrSpc =
       CGM.getContext().getTargetAddressSpace(LangAS::opencl_global);
     PipeTy = llvm::PointerType::get(llvm::StructType::create(
       CGM.getLLVMContext(), "opencl.pipe_t"), PipeAddrSpc);
   }
 
   return PipeTy;
 }
+
+llvm::PointerType *CGOpenCLRuntime::getSamplerType() {
+  if (!SamplerTy)
+    SamplerTy = llvm::PointerType::get(llvm::StructType::create(
+      CGM.getLLVMContext(), "__opencl_sampler_t"),
+      CGM.getContext().getTargetAddressSpace(
+      LangAS::opencl_constant));
+  return SamplerTy;
+}

lib/CodeGen/CodeGenModule.h

   void AddVTableTypeMetadata(llvm::GlobalVariable *VTable, CharUnits Offset,
                              const CXXRecordDecl *RD);
 
   /// \breif Get the declaration of std::terminate for the platform.
   llvm::Constant *getTerminateFn();
 
   llvm::SanitizerStatReport &getSanStats();
 
+  llvm::Value *
+  createOpenCLIntToSamplerConversion(const Expr *E, CodeGenFunction &CGF);
+
 private:
   llvm::Constant *
   GetOrCreateLLVMFunction(StringRef MangledName, llvm::Type *Ty, GlobalDecl D,
                           bool ForVTable, bool DontDefer = false,
                           bool IsThunk = false,
                           llvm::AttributeSet ExtraAttrs = llvm::AttributeSet(),
                           bool IsForDefinition = false);
 

lib/CodeGen/CodeGenModule.cpp

   QualType ASTTy = D->getType();
   CXXRecordDecl *RD = ASTTy->getBaseElementTypeUnsafe()->getAsCXXRecordDecl();
   bool NeedsGlobalCtor = false;
   bool NeedsGlobalDtor = RD && !RD->hasTrivialDestructor();
 
   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.
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
       D->hasAttr<CUDASharedAttr>())
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   else if (!InitExpr) {
     // This is a tentative definition; tentative definitions are

 
     // Now populate the feature map, first with the TargetCPU which is either
     // the default or a new one from the target attribute string. Then we'll use
     // the passed in features (FeaturesAsWritten) along with the new ones from
     // the attribute.
     Target.initFeatureMap(FeatureMap, getDiags(), TargetCPU, ParsedAttr.first);
   } else {
     Target.initFeatureMap(FeatureMap, getDiags(), TargetCPU,
                           Target.getTargetOpts().Features);
   }
 }
 
 llvm::SanitizerStatReport &CodeGenModule::getSanStats() {
   if (!SanStats)
     SanStats = llvm::make_unique<llvm::SanitizerStatReport>(&getModule());
 
   return *SanStats;
 }
+llvm::Value *
+CodeGenModule::createOpenCLIntToSamplerConversion(const Expr *E,
+                                                  CodeGenFunction &CGF) {
+  llvm::Constant *C = EmitConstantExpr(E, E->getType(), &CGF);
+  auto SamplerT = getOpenCLRuntime().getSamplerType();
+  auto FTy = llvm::FunctionType::get(SamplerT, {C->getType()}, false);
+  return CGF.Builder.CreateCall(CreateRuntimeFunction(FTy,
+                                "__translate_sampler_initializer"),
+                                {C});
+}

lib/Edit/RewriteObjCFoundationAPI.cpp

     case CK_ARCProduceObject:
     case CK_ARCConsumeObject:
     case CK_ARCReclaimReturnedObject:
     case CK_ARCExtendBlockObject:
     case CK_NonAtomicToAtomic:
     case CK_CopyAndAutoreleaseBlockObject:
     case CK_BuiltinFnToFnPtr:
     case CK_ZeroToOCLEvent:
+    case CK_IntToOCLSampler:
       return false;
 
     case CK_BooleanToSignedIntegral:
       llvm_unreachable("OpenCL-specific cast in Objective-C?");
     }
   }
 
   if (needsCast) {

lib/Frontend/CompilerInvocation.cpp

   // FIXME: ParsePreprocessorArgs uses the FileManager to read the contents of
   // PCH file and find the original header name. Remove the need to do that in
   // ParsePreprocessorArgs and remove the FileManager
   // parameters from the function and the "FileManager.h" #include.
   FileManager FileMgr(Res.getFileSystemOpts());
   ParsePreprocessorArgs(Res.getPreprocessorOpts(), Args, FileMgr, Diags);
   ParsePreprocessorOutputArgs(Res.getPreprocessorOutputOpts(), Args,
                               Res.getFrontendOpts().ProgramAction);
+
+  // Turn on -Wspir-compat for SPIR target.
+  llvm::Triple T(Res.getTargetOpts().Triple);
+  auto Arch = T.getArch();
+  if (Arch == llvm::Triple::spir || Arch == llvm::Triple::spir64) {
+    Res.getDiagnosticOpts().Warnings.push_back("spir-compat");
+  }
   return Success;
 }
 
 std::string CompilerInvocation::getModuleHash() const {
   // Note: For QoI reasons, the things we use as a hash here should all be
   // dumped via the -module-info flag.
   using llvm::hash_code;
   using llvm::hash_value;

lib/Sema/SemaExpr.cpp

   // struct A -> struct B
   if (isa<TagType>(LHSType) && isa<TagType>(RHSType)) {
     if (Context.typesAreCompatible(LHSType, RHSType)) {
       Kind = CK_NoOp;
       return Compatible;
     }
   }
 
+  if (LHSType->isSamplerT() && RHSType->isIntegerType()) {
+    Kind = CK_IntToOCLSampler;
+    return Compatible;
+  }
+
   return Incompatible;
 }
 
 /// \brief Constructs a transparent union from an expression that is
 /// used to initialize the transparent union.
 static void ConstructTransparentUnion(Sema &S, ASTContext &C,
                                       ExprResult &EResult, QualType UnionType,
                                       FieldDecl *Field) {

lib/Sema/SemaInit.cpp

   return true;
 }
 
 static bool TryOCLSamplerInitialization(Sema &S,
                                         InitializationSequence &Sequence,
                                         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);
   return true;
 }
 
 //
 // OpenCL 1.2 spec, s6.12.10

 
       // Bind the result, in case the library has given initializer_list a
       // non-trivial destructor.
       if (shouldBindAsTemporary(Entity))
         CurInit = S.MaybeBindToTemporary(CurInit.get());
       break;
     }
 
     case SK_OCLSamplerInit: {
-      assert(Step->Type->isSamplerT() && 
+      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!");
+          break;
+        } else if (const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(Init)) {
+          auto Var = cast<VarDecl>(DRE->getDecl());
+          if (!Var->hasGlobalStorage()) {
+            CurInit = ImplicitCastExpr::Create(S.Context, Step->Type,
+                                               CK_LValueToRValue, CurInit.get(),
+                                               /*BasePath=*/nullptr, VK_RValue);
+            break;
+          }
+          Init = cast<ImplicitCastExpr>(const_cast<Expr*>(
+            Var->getInit()))->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;
+
+      llvm::APSInt Result;
+      Init->EvaluateAsInt(Result, S.Context);
+      const uint64_t SamplerValue = Result.getLimitedValue();
+      // 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)
+        S.Diag(Kind.getLocation(), diag::warn_sampler_initializer_invalid_bits)
+          << "Filter Mode";
+      if (AddressingMode > 4)
+        S.Diag(Kind.getLocation(), diag::warn_sampler_initializer_invalid_bits)
+          << "Addressing Mode";
+
+      CurInit = S.ImpCastExprToType(Init, S.Context.OCLSamplerTy,
+                                      CK_IntToOCLSampler);
       break;
     }
     case SK_OCLZeroEvent: {
       assert(Step->Type->isEventT() && 
              "Event initialization on non-event type.");
 
       CurInit = S.ImpCastExprToType(CurInit.get(), Step->Type,
                                     CK_ZeroToOCLEvent,

lib/StaticAnalyzer/Core/ExprEngineC.cpp

switch (CastE->getCastKind()) {
       case CK_IntegralComplexToBoolean:
       case CK_IntegralComplexCast:
       case CK_IntegralComplexToFloatingComplex:
       case CK_CPointerToObjCPointerCast:
       case CK_BlockPointerToObjCPointerCast:
       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);
         V = svalBuilder.evalCast(V, T, ExTy);
         // Negate the result if we're treating the boolean as a signed i1
         if (CastE->getCastKind() == CK_BooleanToSignedIntegral)
           V = evalMinus(V);
         state = state->BindExpr(CastE, LCtx, V);

test/CodeGenOpenCL/opencl_types.cl

-// 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(%__opencl_sampler_t 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 %__opencl_sampler_t addrspace(2)*
   event_t evt;
   // CHECK: alloca %opencl.event_t*
-  // CHECK: store i32 5,
+  // CHECK: store %__opencl_sampler_t addrspace(2)*
   fnc4smp(smp);
-  // CHECK: call {{.*}}void @fnc4smp(i32
+  // CHECK: call {{.*}}void @fnc4smp(%__opencl_sampler_t addrspace(2)*
   fnc4smp(glb_smp);
-  // CHECK: call {{.*}}void @fnc4smp(i32
+  // CHECK: call {{.*}}void @fnc4smp(%__opencl_sampler_t addrspace(2)*
 }
 
 void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
 // CHECK-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}

test/CodeGenOpenCL/sampler.cl

+// 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
+
+// CHECK: %__opencl_sampler_t = type opaque
+
+constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
+// CHECK-NOT: glb_smp
+
+void fnc4smp(sampler_t s) {}
+// CHECK: define spir_func void @fnc4smp(%__opencl_sampler_t 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 %__opencl_sampler_t addrspace(2)*
+  // CHECK: [[SAMP:%[0-9]+]] = call %__opencl_sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
+  // CHECK: store %__opencl_sampler_t addrspace(2)* [[SAMP]], %__opencl_sampler_t addrspace(2)** [[smp_ptr]]
+
+  fnc4smp(smp);
+  // CHECK-NOT: call %__opencl_sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
+  // CHECK: [[SAMP:%[0-9]+]] = load %__opencl_sampler_t addrspace(2)*, %__opencl_sampler_t addrspace(2)** [[smp_ptr]]
+  // CHECK: call spir_func void @fnc4smp(%__opencl_sampler_t addrspace(2)* [[SAMP]])
+
+  fnc4smp(smp);
+  // CHECK-NOT: call %__opencl_sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
+  // CHECK: [[SAMP:%[0-9]+]] = load %__opencl_sampler_t addrspace(2)*, %__opencl_sampler_t addrspace(2)** [[smp_ptr]]
+  // CHECK: call spir_func void @fnc4smp(%__opencl_sampler_t addrspace(2)* [[SAMP]])
+
+  fnc4smp(glb_smp);
+  // CHECK: [[SAMP:%[0-9]+]] = call %__opencl_sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
+  // CHECK: call spir_func void @fnc4smp(%__opencl_sampler_t addrspace(2)* [[SAMP]])
+
+}

test/SemaOpenCL/sampler_t.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -DCHECK_SAMPLER_VALUE -Wspir-compat -triple amdgcn--amdhsa
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -DCHECK_SAMPLER_VALUE -triple spir-unknown-unknown
 
-constant sampler_t glb_smp = 5;
+#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; // expected-error{{variable in constant address space must be initialized}}
+global sampler_t glb_smp3 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST; // expected-error{{sampler type cannot be used with the __local and __global address space qualifiers}}
+
+constant sampler_t glb_smp4 = 0;
+#ifdef CHECK_SAMPLER_VALUE
+// expected-warning@-2{{sampler initializer has invalid Filter Mode bits}}
+#endif
+
+constant sampler_t glb_smp5 = 0x1f;
+#ifdef CHECK_SAMPLER_VALUE
+// expected-warning@-2{{sampler initializer has invalid Addressing Mode bits}}
+#endif
 
 void foo(sampler_t);
 
 constant struct sampler_s {
   sampler_t smp; // expected-error{{the 'sampler_t' type cannot be used to declare a structure or union field}}
 } sampler_str = {0};
 
 void kernel ker(sampler_t argsmp) {
   local sampler_t smp; // expected-error{{sampler type cannot be used with the __local and __global address space qualifiers}}
-  const sampler_t const_smp = 7;
+  const sampler_t const_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
+  const sampler_t const_smp2;
   foo(glb_smp);
+  foo(glb_smp2);
+  foo(glb_smp3);
   foo(const_smp);
+  foo(const_smp2);
+  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}}
 }
 
 void bad(sampler_t*); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
 
 void bar() {
-  sampler_t smp1 = 7;
-  sampler_t smp2 = 2;
+  sampler_t smp1 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
+  sampler_t smp2 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST;
   smp1=smp2; //expected-error{{invalid operands to binary expression ('sampler_t' and 'sampler_t')}}
   smp1+1; //expected-error{{invalid operands to binary expression ('sampler_t' and 'int')}}
   &smp1; //expected-error{{invalid argument type 'sampler_t' to unary expression}}
   *smp2; //expected-error{{invalid argument type 'sampler_t' to unary expression}}
 }
 
-sampler_t bad(); //expected-error{{declaring function return value of type 'sampler_t' is not allowed}}
+sampler_t bad(void); //expected-error{{declaring function return value of type 'sampler_t' is not allowed}}