diff --git a/include/LLVMSPIRVOpts.h b/include/LLVMSPIRVOpts.h
index c956eb7ab4..bd1c73f318 100644
--- a/include/LLVMSPIRVOpts.h
+++ b/include/LLVMSPIRVOpts.h
@@ -221,6 +221,14 @@ class TranslatorOpts {
     PreserveOCLKernelArgTypeMetadataThroughString = Value;
   }
 
+  bool shouldEmitFunctionPtrAddrSpace() const noexcept {
+    return EmitFunctionPtrAddrSpace;
+  }
+
+  void setEmitFunctionPtrAddrSpace(bool Value) noexcept {
+    EmitFunctionPtrAddrSpace = Value;
+  }
+
 private:
   // Common translation options
   VersionNumber MaxVersion = VersionNumber::MaximumVersion;
@@ -262,6 +270,10 @@ class TranslatorOpts {
   // kernel_arg_type_qual metadata through OpString
   bool PreserveOCLKernelArgTypeMetadataThroughString = false;
 
+  // Controls if CodeSectionINTEL can be emitted and consumed with a dedicated
+  // address space
+  bool EmitFunctionPtrAddrSpace = false;
+
   bool PreserveAuxData = false;
 };
 
diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h
index f0acfce59c..4c78d27b66 100644
--- a/lib/SPIRV/SPIRVInternal.h
+++ b/lib/SPIRV/SPIRVInternal.h
@@ -193,6 +193,7 @@ enum SPIRAddressSpace {
   SPIRAS_GlobalHost,
   SPIRAS_Input,
   SPIRAS_Output,
+  SPIRAS_CodeSectionINTEL,
   SPIRAS_Count,
 };
 
@@ -203,6 +204,8 @@ template <> inline void SPIRVMap<SPIRAddressSpace, std::string>::init() {
   add(SPIRAS_Local, "Local");
   add(SPIRAS_Generic, "Generic");
   add(SPIRAS_Input, "Input");
+  add(SPIRAS_CodeSectionINTEL, "CodeSectionINTEL");
+
   add(SPIRAS_GlobalDevice, "GlobalDevice");
   add(SPIRAS_GlobalHost, "GlobalHost");
 }
@@ -219,6 +222,7 @@ inline void SPIRVMap<SPIRAddressSpace, SPIRVStorageClassKind>::init() {
   add(SPIRAS_Input, StorageClassInput);
   add(SPIRAS_GlobalDevice, StorageClassDeviceOnlyINTEL);
   add(SPIRAS_GlobalHost, StorageClassHostOnlyINTEL);
+  add(SPIRAS_CodeSectionINTEL, StorageClassCodeSectionINTEL);
 }
 typedef SPIRVMap<SPIRAddressSpace, SPIRVStorageClassKind> SPIRSPIRVAddrSpaceMap;
 
diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp
index 011ac07974..47315f77cb 100644
--- a/lib/SPIRV/SPIRVReader.cpp
+++ b/lib/SPIRV/SPIRVReader.cpp
@@ -381,11 +381,17 @@ Type *SPIRVToLLVM::transType(SPIRVType *T, bool IsClassMember) {
   }
   case internal::OpTypeTokenINTEL:
     return mapType(T, Type::getTokenTy(*Context));
-  case OpTypePointer:
+  case OpTypePointer: {
+    unsigned AS = SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass());
+    if (AS == SPIRAS_CodeSectionINTEL && !BM->shouldEmitFunctionPtrAddrSpace())
+      AS = SPIRAS_Private;
+    if (BM->shouldEmitFunctionPtrAddrSpace() &&
+        T->getPointerElementType()->getOpCode() == OpTypeFunction)
+      AS = SPIRAS_CodeSectionINTEL;
     return mapType(
         T, PointerType::get(
-               transType(T->getPointerElementType(), IsClassMember),
-               SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass())));
+               transType(T->getPointerElementType(), IsClassMember), AS));
+  }
   case OpTypeVector:
     return mapType(T,
                    FixedVectorType::get(transType(T->getVectorComponentType()),
@@ -1564,10 +1570,20 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F,
     case OpTypeVector:
       return mapValue(BV, ConstantVector::get(CV));
     case OpTypeMatrix:
-    case OpTypeArray:
-      return mapValue(
-          BV, ConstantArray::get(dyn_cast<ArrayType>(transType(BCC->getType())),
-                                 CV));
+    case OpTypeArray: {
+      auto *AT = cast<ArrayType>(transType(BCC->getType()));
+      for (size_t I = 0; I != AT->getNumElements(); ++I) {
+        auto *ElemTy = AT->getElementType();
+        if (auto *ElemPtrTy = dyn_cast<PointerType>(ElemTy)) {
+          assert(isa<PointerType>(CV[I]->getType()) &&
+                 "Constant type doesn't match constexpr array element type");
+          if (ElemPtrTy->getAddressSpace() !=
+              cast<PointerType>(CV[I]->getType())->getAddressSpace())
+            CV[I] = ConstantExpr::getAddrSpaceCast(CV[I], AT->getElementType());
+        }
+      }
+      return mapValue(BV, ConstantArray::get(AT, CV));
+    }
     case OpTypeStruct: {
       auto BCCTy = dyn_cast<StructType>(transType(BCC->getType()));
       auto Members = BCCTy->getNumElements();
@@ -1582,7 +1598,12 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F,
               !BCCTy->getElementType(I)->isPointerTy())
             continue;
 
-          CV[I] = ConstantExpr::getBitCast(CV[I], BCCTy->getElementType(I));
+          if (cast<PointerType>(CV[I]->getType())->getAddressSpace() !=
+              cast<PointerType>(BCCTy->getElementType(I))->getAddressSpace())
+            CV[I] =
+                ConstantExpr::getAddrSpaceCast(CV[I], BCCTy->getElementType(I));
+          else
+            CV[I] = ConstantExpr::getBitCast(CV[I], BCCTy->getElementType(I));
         }
       }
 
@@ -1620,7 +1641,10 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F,
         static_cast<SPIRVConstantFunctionPointerINTEL *>(BV);
     SPIRVFunction *F = BC->getFunction();
     BV->setName(F->getName());
-    return mapValue(BV, transFunction(F));
+    const unsigned AS = BM->shouldEmitFunctionPtrAddrSpace()
+                            ? SPIRAS_CodeSectionINTEL
+                            : SPIRAS_Private;
+    return mapValue(BV, transFunction(F, AS));
   }
 
   case OpUndef:
@@ -3038,7 +3062,7 @@ void SPIRVToLLVM::transFunctionAttrs(SPIRVFunction *BF, Function *F) {
   });
 }
 
-Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) {
+Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF, unsigned AS) {
   auto Loc = FuncMap.find(BF);
   if (Loc != FuncMap.end())
     return Loc->second;
@@ -3087,7 +3111,7 @@ Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) {
   }
   Function *F = M->getFunction(FuncName);
   if (!F)
-    F = Function::Create(FT, Linkage, FuncName, M);
+    F = Function::Create(FT, Linkage, AS, FuncName, M);
   F = cast<Function>(mapValue(BF, F));
   mapFunction(BF, F);
 
@@ -3486,6 +3510,17 @@ bool SPIRVToLLVM::translate() {
     DbgTran->transDebugInst(EI);
   }
 
+  for (auto *FP : BM->getFunctionPointers()) {
+    SPIRVConstantFunctionPointerINTEL *BC =
+        static_cast<SPIRVConstantFunctionPointerINTEL *>(FP);
+    SPIRVFunction *F = BC->getFunction();
+    FP->setName(F->getName());
+    const unsigned AS = BM->shouldEmitFunctionPtrAddrSpace()
+                            ? SPIRAS_CodeSectionINTEL
+                            : SPIRAS_Private;
+    mapValue(FP, transFunction(F, AS));
+  }
+
   for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) {
     transFunction(BM->getFunction(I));
     transUserSemantic(BM->getFunction(I));
diff --git a/lib/SPIRV/SPIRVReader.h b/lib/SPIRV/SPIRVReader.h
index f1de7c6f52..ef4b6c3e1b 100644
--- a/lib/SPIRV/SPIRVReader.h
+++ b/lib/SPIRV/SPIRVReader.h
@@ -98,7 +98,7 @@ class SPIRVToLLVM {
   void transAuxDataInst(SPIRVExtInst *BC);
   std::vector<Value *> transValue(const std::vector<SPIRVValue *> &,
                                   Function *F, BasicBlock *);
-  Function *transFunction(SPIRVFunction *F);
+  Function *transFunction(SPIRVFunction *F, unsigned AS = SPIRAS_Private);
   void transFunctionAttrs(SPIRVFunction *BF, Function *F);
   Value *transBlockInvoke(SPIRVValue *Invoke, BasicBlock *BB);
   Instruction *transWGSizeQueryBI(SPIRVInstruction *BI, BasicBlock *BB);
diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp
index 54af00790d..c51ab56d9e 100644
--- a/lib/SPIRV/SPIRVWriter.cpp
+++ b/lib/SPIRV/SPIRVWriter.cpp
@@ -508,6 +508,11 @@ SPIRVType *LLVMToSPIRVBase::transPointerType(Type *ET, unsigned AddrSpc) {
       ((AddrSpc == SPIRAS_GlobalDevice) || (AddrSpc == SPIRAS_GlobalHost))) {
     return transPointerType(ET, SPIRAS_Global);
   }
+  // Lower function pointer address space to private if
+  // spirv-emit-function-ptr-addr-space is not passed
+  if (AddrSpc == SPIRAS_CodeSectionINTEL &&
+      !BM->shouldEmitFunctionPtrAddrSpace())
+    return transPointerType(ET, SPIRAS_Private);
   if (ST && !ST->isSized()) {
     Op OpCode;
     StringRef STName = ST->getName();
@@ -615,6 +620,10 @@ SPIRVType *LLVMToSPIRVBase::transPointerType(SPIRVType *ET, unsigned AddrSpc) {
   if (Loc != PointeeTypeMap.end())
     return Loc->second;
 
+  if (AddrSpc == SPIRAS_CodeSectionINTEL &&
+      !BM->shouldEmitFunctionPtrAddrSpace())
+    return transPointerType(ET, SPIRAS_Private);
+
   SPIRVType *TranslatedTy = BM->addPointerType(
       SPIRSPIRVAddrSpaceMap::map(static_cast<SPIRAddressSpace>(AddrSpc)), ET);
   PointeeTypeMap[TypeKey] = TranslatedTy;
diff --git a/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp b/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp
index 4612b4275c..61ad987baf 100644
--- a/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp
+++ b/lib/SPIRV/libSPIRV/SPIRVInstruction.cpp
@@ -286,6 +286,11 @@ SPIRVInstruction *createInstFromSpecConstantOp(SPIRVSpecConstantOp *Inst) {
   auto OC = static_cast<Op>(Ops[0]);
   assert(isSpecConstantOpAllowedOp(OC) &&
          "Op code not allowed for OpSpecConstantOp");
+  auto *Const = Inst->getOperand(1);
+  // LLVM would eliminate a bitcast from a function pointer in a constexpr
+  // context. Cut this short here to avoid necessity to align address spaces
+  if (OC == OpBitcast && Const->getOpCode() == OpConstantFunctionPointerINTEL)
+    return static_cast<SPIRVInstruction *>(Const);
   Ops.erase(Ops.begin(), Ops.begin() + 1);
   auto *BM = Inst->getModule();
   auto *RetInst = SPIRVInstTemplateBase::create(
diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/lib/SPIRV/libSPIRV/SPIRVModule.cpp
index 6829200d7c..1fc6ef7461 100644
--- a/lib/SPIRV/libSPIRV/SPIRVModule.cpp
+++ b/lib/SPIRV/libSPIRV/SPIRVModule.cpp
@@ -168,6 +168,13 @@ class SPIRVModuleImpl : public SPIRVModule {
   SPIRVConstant *getLiteralAsConstant(unsigned Literal) override;
   unsigned getNumFunctions() const override { return FuncVec.size(); }
   unsigned getNumVariables() const override { return VariableVec.size(); }
+  std::vector<SPIRVValue *> getFunctionPointers() const override {
+    std::vector<SPIRVValue *> Res;
+    for (auto *C : ConstVec)
+      if (C->getOpCode() == OpConstantFunctionPointerINTEL)
+        Res.emplace_back(C);
+    return Res;
+  }
   SourceLanguage getSourceLanguage(SPIRVWord *Ver = nullptr) const override {
     if (Ver)
       *Ver = SrcLangVer;
diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.h b/lib/SPIRV/libSPIRV/SPIRVModule.h
index 6e69644473..41e18e5b31 100644
--- a/lib/SPIRV/libSPIRV/SPIRVModule.h
+++ b/lib/SPIRV/libSPIRV/SPIRVModule.h
@@ -140,6 +140,7 @@ class SPIRVModule {
   virtual SPIRVMemoryModelKind getMemoryModel() const = 0;
   virtual unsigned getNumFunctions() const = 0;
   virtual unsigned getNumVariables() const = 0;
+  virtual std::vector<SPIRVValue *> getFunctionPointers() const = 0;
   virtual SourceLanguage getSourceLanguage(SPIRVWord *) const = 0;
   virtual std::set<std::string> &getSourceExtension() = 0;
   virtual SPIRVValue *getValue(SPIRVId TheId) const = 0;
@@ -539,6 +540,10 @@ class SPIRVModule {
         .shouldPreserveOCLKernelArgTypeMetadataThroughString();
   }
 
+  bool shouldEmitFunctionPtrAddrSpace() const noexcept {
+    return TranslationOpts.shouldEmitFunctionPtrAddrSpace();
+  }
+
   bool preserveAuxData() const noexcept {
     return TranslationOpts.preserveAuxData();
   }
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll
new file mode 100644
index 0000000000..7293f48c71
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/bitcast.ll
@@ -0,0 +1,51 @@
+; OpenCL C source:
+; char foo(char a) {
+;     return a;
+; }
+; void bar() {
+;    int (*fun_ptr)(int) = &foo;
+; }
+
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv %t.spv -to-text -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+
+; CHECK-SPIRV: TypeFunction [[#FOO_TY:]] [[#]] [[#]]
+; CHECK-SPIRV: TypeFunction [[#DEST_TY:]] [[#]] [[#]]
+; CHECK-SPIRV: TypePointer [[#DEST_TY_PTR:]] [[#]] [[#DEST_TY]]
+; CHECK-SPIRV: TypePointer [[#FOO_TY_PTR:]] [[#]] [[#FOO_TY]]
+; CHECK-SPIRV: ConstantFunctionPointerINTEL [[#FOO_TY_PTR]] [[#FOO_PTR:]] [[#FOO:]]
+; CHECK-SPIRV: Function [[#]] [[#FOO]] [[#]] [[#FOO_TY]]
+
+; CHECK-SPIRV: Bitcast [[#DEST_TY_PTR]] [[#]] [[#FOO_PTR]]
+
+; CHECK-LLVM: bitcast i8 (i8) addrspace(9)* @foo to i32 (i32) addrspace(9)*
+
+; ModuleID = './example.c'
+source_filename = "./example.c"
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir"
+
+; Function Attrs: noinline nounwind optnone
+define dso_local spir_func signext i8 @foo(i8 signext %0) #0 {
+  ret i8 %0
+}
+
+; Function Attrs: noinline nounwind optnone
+define dso_local spir_func void @bar() #0 {
+  %1 = alloca i32 (i32)*, align 4
+  store i32 (i32)* bitcast (i8 (i8)* @foo to i32 (i32)*), i32 (i32)** %1, align 4
+  ret void
+}
+
+attributes #0 = { noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+
+!llvm.module.flags = !{!0}
+!llvm.ident = !{!1}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{!"clang version 11.0.0 (https://github.com/llvm/llvm-project.git 0e1accd0f726eef2c47be9f37dd0a06cb50d207e)"}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll
new file mode 100644
index 0000000000..c937ae11b3
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/const-function-pointer.ll
@@ -0,0 +1,67 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+; CHECK-SPIRV: Name [[F1Name:[0-9]+]] "f1"
+; CHECK-SPIRV: Name [[F2Name:[0-9]+]] "f2"
+; CHECK-SPIRV: TypeInt [[Int32:[0-9]+]] 32
+; CHECK-SPIRV: TypeInt [[Int64:[0-9]+]] 64
+; CHECK-SPIRV-DAG: Constant [[Int32]] [[XArg:[0-9]+]] 32
+; CHECK-SPIRV-DAG: Constant [[Int32]] [[YArg:[0-9]+]] 2
+
+; CHECK-SPIRV: ConstantFunctionPointerINTEL {{[0-9]+}} [[F1:[0-9]+]] [[F1Name]]
+; CHECK-SPIRV: ConstantFunctionPointerINTEL {{[0-9]+}} [[F2:[0-9]+]] [[F2Name]]
+; CHECK-SPIRV: ConstantComposite {{[0-9]+}} [[ConstComp:[0-9]+]] [[F1]] [[F2]]
+; CHECK-SPIRV: Variable {{[0-9]+}} [[Var:[0-9]+]] {{[0-9]+}} [[ConstComp]]
+
+; CHECK-SPIRV: InBoundsPtrAccessChain {{[0-9]+}} [[GEP:[0-9]+]] [[Var]] {{[0-9]+}} {{[0-9]+}}
+; CHECK-SPIRV: Load {{[0-9]+}} [[FuncPtr:[0-9]+]] [[GEP]]
+; CHECK-SPIRV: FunctionPointerCallINTEL [[Int32]] {{[0-9]+}} [[FuncPtr]] [[XArg]] [[YArg]]
+
+; CHECK-LLVM: @__const.main.funcs = internal constant [2 x i32 (i32, i32) addrspace(9)*] [i32 (i32, i32) addrspace(9)* @f1, i32 (i32, i32) addrspace(9)* @f2], align 16
+; CHECK-LLVM:   %[[Idx:[a-z0-9]+]] = getelementptr inbounds [2 x i32 (i32, i32) addrspace(9)*], [2 x i32 (i32, i32) addrspace(9)*]* @__const.main.funcs, i64 0, i64 %{{[a-z0-9]+}}
+; CHECK-LLVM:   %[[FuncPtr:[a-z0-9]+]] = load i32 (i32, i32) addrspace(9)*, i32 (i32, i32) addrspace(9)** %[[Idx]], align 8
+; CHECK-LLVM:   %{{[a-z0-9]+}} = call spir_func addrspace(9) i32 %[[FuncPtr]](i32 32, i32 2)
+
+target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "spir-unknown-unknown"
+
+@__const.main.funcs = private unnamed_addr constant [2 x i32 (i32, i32)*] [i32 (i32, i32)* @f1, i32 (i32, i32)* @f2], align 16
+
+; Function Attrs: norecurse nounwind readnone uwtable
+define dso_local i32 @f1(i32 %a, i32 %b) #0 {
+entry:
+  %add = add nsw i32 %b, %a
+  ret i32 %add
+}
+
+; Function Attrs: norecurse nounwind readnone uwtable
+define dso_local i32 @f2(i32 %a, i32 %b) #0 {
+entry:
+  %sub = sub nsw i32 %a, %b
+  ret i32 %sub
+}
+
+; Function Attrs: nounwind uwtable
+define dso_local i32 @main() local_unnamed_addr #1 {
+entry:
+  %call = tail call i32 @rand() #3
+  %rem = srem i32 %call, 2
+  %idxprom = sext i32 %rem to i64
+  %arrayidx = getelementptr inbounds [2 x i32 (i32, i32)*], [2 x i32 (i32, i32)*]* @__const.main.funcs, i64 0, i64 %idxprom
+  %0 = load i32 (i32, i32)*, i32 (i32, i32)** %arrayidx, align 8
+  %call1 = tail call i32 %0(i32 32, i32 2) #3
+  ret i32 %call1
+}
+
+; Function Attrs: nounwind
+declare dso_local i32 @rand() local_unnamed_addr #2
+
+attributes #0 = { norecurse nounwind readnone uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #2 = { nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll
new file mode 100644
index 0000000000..edb8d910c2
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll
@@ -0,0 +1,67 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o %t.spt -spirv-text -spirv-ext=+SPV_INTEL_function_pointers
+; RUN: FileCheck < %t.spt %s --check-prefix CHECK-SPIRV
+
+; RUN: llvm-spirv %t.spt -o %t.spv -to-binary
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o %t.rev.ll
+; RUN: FileCheck < %t.rev.ll %s --check-prefix CHECK-LLVM
+
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+
+; CHECK-SPIRV: Decorate [[#TargetId:]] ArgumentAttributeINTEL 0 4
+; CHECK-SPIRV: Decorate [[#TargetId]] ArgumentAttributeINTEL 0 5
+; CHECK-SPIRV: Decorate [[#TargetId]] ArgumentAttributeINTEL 0 2
+; CHECK-SPIRV: FunctionPointerCallINTEL
+; CHECK-SPIRV-SAME: [[#TargetId]]
+
+; CHECK-LLVM: call spir_func addrspace(9) void %cond.i.i(%multi_ptr* noalias nocapture byval(%multi_ptr) %agg.tmp.i.i)
+
+; ModuleID = 'sycl_test.cpp'
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"multi_ptr" = type { i32* }
+%"range" = type { %"array" }
+%"array" = type { [1 x i64] }
+%wrapper_class = type { i32 addrspace(1)* }
+%wrapper_class.0 = type { i32 addrspace(1)* }
+
+$RoundedRangeKernel = comdat any
+
+; Function Attrs: nounwind
+define spir_func void @inc_function(%"multi_ptr"* byval(%"multi_ptr") noalias nocapture %ptr) #0 {
+entry:
+  ret void
+}
+
+
+; Function Attrs: convergent norecurse
+define weak_odr dso_local spir_kernel void @RoundedRangeKernel(%"range"* byval(%"range") align 8 %_arg_NumWorkItems, i1 zeroext %_arg_, %wrapper_class* byval(%wrapper_class) align 8 %_arg_1, %wrapper_class.0* byval(%wrapper_class.0) align 8 %_arg_2) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !6 {
+entry:
+  %agg.tmp.i.i = alloca %"multi_ptr", align 8
+  %cond.i.i = select i1 %_arg_, void (%"multi_ptr"*)* @inc_function, void (%"multi_ptr"*)* null
+  call spir_func void %cond.i.i(%"multi_ptr"* nonnull byval(%"multi_ptr") align 8 noalias nocapture %agg.tmp.i.i) #1, !callees !7
+  ret void
+}
+
+attributes #0 = { convergent norecurse "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="sycl_test.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="true" }
+attributes #1 = { convergent }
+
+!llvm.module.flags = !{!0, !1}
+!opencl.spir.version = !{!2}
+!spirv.Source = !{!3}
+!opencl.used.extensions = !{!4}
+!opencl.used.optional.core.features = !{!4}
+!opencl.compiler.options = !{!4}
+!llvm.ident = !{!5}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"frame-pointer", i32 2}
+!2 = !{i32 1, i32 2}
+!3 = !{i32 4, i32 100000}
+!4 = !{}
+!5 = !{!"Compiler"}
+!6 = !{i32 -1, i32 -1, i32 -1, i32 -1}
+!7 = !{void (%"multi_ptr"*)* @inc_function}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll
new file mode 100644
index 0000000000..0c1394d48b
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/fp-from-host.ll
@@ -0,0 +1,70 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+;
+; Generated from:
+; typedef int (*fp_t)(int);
+;
+; __kernel void test(__global int *fp, __global int *data) {
+;
+;   data[0] = ((fp_t)(*fp))(data[1]);
+; }
+;
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+;
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: TypeInt [[INT32_TYPE_ID:[0-9]+]] 32
+; CHECK-SPIRV: TypePointer [[INT_PTR:[0-9]+]] 5 [[INT32_TYPE_ID]]
+; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[INT32_TYPE_ID]] [[INT32_TYPE_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_TYPE_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]]
+; CHECK-SPIRV: FunctionParameter [[INT_PTR]] [[FP:[0-9]+]]
+; CHECK-SPIRV: Load [[INT32_TYPE_ID]] [[FUNC_ADDR:[0-9]+]] [[FP]]
+; CHECK-SPIRV: ConvertUToPtr [[FOO_TYPE_PTR_ID]] [[FOO_PTR:[0-9]+]] [[FUNC_ADDR]]
+; CHECK-SPIRV: FunctionPointerCallINTEL [[INT32_TYPE_ID]] {{[0-9]+}} [[FOO_PTR]]
+;
+; CHECK-LLVM: define spir_kernel void @test(i32 addrspace(1)*
+; CHECK-LLVM: %{{.*}} = call spir_func addrspace(9) i32 %{{.*}}(i32
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; Function Attrs: convergent nounwind
+define spir_kernel void @test(i32 addrspace(1)* %fp, i32 addrspace(1)* %data) #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
+entry:
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %data, i64 1
+  %0 = load i32, i32 addrspace(1)* %arrayidx, align 4, !tbaa !8
+  %1 = load i32, i32 addrspace(1)* %fp, align 4, !tbaa !8
+  %2 = inttoptr i32 %1 to i32 (i32)*
+  %call = call spir_func i32 %2(i32 %0) #1
+  %arrayidx1 = getelementptr inbounds i32, i32 addrspace(1)* %data, i64 0
+  store i32 %call, i32 addrspace(1)* %arrayidx1, align 4, !tbaa !8
+  ret void
+}
+
+attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { convergent }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 0}
+!2 = !{i32 1, i32 2}
+!3 = !{!"clang version 7.1.0 "}
+!4 = !{i32 1, i32 1}
+!5 = !{!"none", !"none"}
+!6 = !{!"int*", !"int*"}
+!7 = !{!"", !""}
+!8 = !{!9, !9, i64 0}
+!9 = !{!"int", !10, i64 0}
+!10 = !{!"omnipotent char", !11, i64 0}
+!11 = !{!"Simple C/C++ TBAA"}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll
new file mode 100644
index 0000000000..7ee0e49daa
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-as-function-arg.ll
@@ -0,0 +1,177 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+;
+; Generated from:
+; int helper(int (*f)(int), int arg) {
+;   return f(arg);
+; }
+;
+; int foo(int v) {
+;   return v + 1;
+; }
+;
+; int bar(int v) {
+;   return v + 2;
+; }
+;
+; __kernel void test(__global int *data, int control) {
+;   int (*fp)(int) = 0;
+;
+;   if (get_global_id(0) % control == 0)
+;     fp = &foo;
+;   else
+;     fp = &bar;
+;
+;   data[get_global_id(0)] = helper(fp, data[get_global_id(0)]);
+; }
+;
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+;
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9]+]] 32
+; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]
+; CHECK-SPIRV: TypeFunction [[HELPER_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[FOO_PTR_TYPE_ID]] [[TYPE_INT32_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_PTR_ALLOCA_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_PTR_TYPE_ID]]
+; CHECK-SPIRV: TypePointer [[TYPE_INT32_ALLOCA_ID:[0-9]+]] {{[0-9]+}} [[TYPE_INT32_ID]]
+; CHECK-SPIRV: FunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[FOO_PTR_ID:[0-9]+]] [[FOO_ID:[0-9]+]]
+; CHECK-SPIRV: FunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[BAR_PTR_ID:[0-9]+]] [[BAR_ID:[0-9]+]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[HELPER_ID:[0-9]+]] {{[0-9]+}} [[HELPER_TYPE_ID]]
+; CHECK-SPIRV: FunctionParameter [[FOO_PTR_TYPE_ID]] [[T_PTR_ARG_ID:[0-9]+]]
+; CHECK-SPIRV: FunctionParameter [[TYPE_INT32_ID:[0-9]+]] [[INT_ARG_ID:[0-9]+]]
+; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_TYPE_ID]] [[T_PTR_ALLOCA_ID:[0-9]+]]
+; CHECK-SPIRV: Variable [[TYPE_INT32_ALLOCA_ID]] [[INT_ALLOCA_ID:[0-9]+]]
+; CHECK-SPIRV: Store [[T_PTR_ALLOCA_ID]] [[T_PTR_ARG_ID]]
+; CHECK-SPIRV: Store [[INT_ALLOCA_ID]] [[INT_ARG_ID]]
+; CHECK-SPIRV: Load [[FOO_PTR_TYPE_ID]] [[LOADED_T_PTR:[0-9]+]] [[T_PTR_ALLOCA_ID]]
+; CHECK-SPIRV: Load [[TYPE_INT32_ID]] [[LOADED_INT:[0-9]+]] [[INT_ALLOCA_ID]]
+; CHECK-SPIRV: FunctionPointerCallINTEL [[TYPE_INT32_ID]] [[RESULT:[0-9]+]] [[LOADED_T_PTR]] [[LOADED_INT]]
+; CHECK-SPIRV: ReturnValue [[RESULT]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] {{[0-9]+}} [[FOO_TYPE_ID]]
+; CHECK-SPIRV: Function {{[0-9]+}} [[BAR_ID]] {{[0-9]+}} [[FOO_TYPE_ID]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]]
+; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_TYPE_ID]] [[F_PTR_ALLOCA_ID:[0-9]+]]
+; CHECK-SPIRV: Store [[F_PTR_ALLOCA_ID]] [[FOO_PTR_ID]]
+; CHECK-SPIRV: Store [[F_PTR_ALLOCA_ID]] [[BAR_PTR_ID]]
+; CHECK-SPIRV: Load [[FOO_PTR_TYPE_ID]] [[LOADED_F_PTR:[0-9]+]] [[F_PTR_ALLOCA_ID]]
+; CHECK-SPIRV: FunctionCall {{[0-9]+}} {{[0-9]+}} [[HELPER_ID]] [[LOADED_F_PTR]]
+;
+; CHECK-LLVM: define spir_func i32 @helper(i32 (i32) addrspace(9)* %[[F:.*]],
+; CHECK-LLVM: %[[F_ADDR:.*]] = alloca i32 (i32) addrspace(9)*
+; CHECK-LLVM: store i32 (i32) addrspace(9)* %[[F]], i32 (i32) addrspace(9)** %[[F_ADDR]]
+; CHECK-LLVM: %[[F_LOADED:.*]] = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %[[F_ADDR]]
+; CHECK-LLVM: %[[CALL:.*]] = call spir_func addrspace(9) i32 %[[F_LOADED]]
+; CHECK-LLVM: ret i32 %[[CALL]]
+;
+; CHECK-LLVM: define spir_kernel void @test
+; CHECK-LLVM: %[[FP:.*]] = alloca i32 (i32) addrspace(9)*
+; CHECK-LLVM: store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %[[FP]]
+; CHECK-LLVM: store i32 (i32) addrspace(9)* @bar, i32 (i32) addrspace(9)** %[[FP]]
+; CHECK-LLVM: %[[FP_LOADED:.*]] = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %[[FP]]
+; CHECK-LLVM: call spir_func i32 @helper(i32 (i32) addrspace(9)* %[[FP_LOADED]]
+
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @helper(i32 (i32)* %f, i32 %arg) #0 {
+entry:
+  %f.addr = alloca i32 (i32)*, align 8
+  %arg.addr = alloca i32, align 4
+  store i32 (i32)* %f, i32 (i32)** %f.addr, align 8
+  store i32 %arg, i32* %arg.addr, align 4
+  %0 = load i32 (i32)*, i32 (i32)** %f.addr, align 8
+  %1 = load i32, i32* %arg.addr, align 4
+  %call = call spir_func i32 %0(i32 %1) #3
+  ret i32 %call
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @foo(i32 %v) #0 {
+entry:
+  %v.addr = alloca i32, align 4
+  store i32 %v, i32* %v.addr, align 4
+  %0 = load i32, i32* %v.addr, align 4
+  %add = add nsw i32 %0, 1
+  ret i32 %add
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @bar(i32 %v) #0 {
+entry:
+  %v.addr = alloca i32, align 4
+  store i32 %v, i32* %v.addr, align 4
+  %0 = load i32, i32* %v.addr, align 4
+  %add = add nsw i32 %0, 2
+  ret i32 %add
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_kernel void @test(i32 addrspace(1)* %data, i32 %control) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 {
+entry:
+  %data.addr = alloca i32 addrspace(1)*, align 8
+  %control.addr = alloca i32, align 4
+  %fp = alloca i32 (i32)*, align 8
+  store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8
+  store i32 %control, i32* %control.addr, align 4
+  store i32 (i32)* null, i32 (i32)** %fp, align 8
+  %call = call spir_func i64 @_Z13get_global_idj(i32 0) #4
+  %0 = load i32, i32* %control.addr, align 4
+  %conv = sext i32 %0 to i64
+  %rem = urem i64 %call, %conv
+  %cmp = icmp eq i64 %rem, 0
+  br i1 %cmp, label %if.then, label %if.else
+
+if.then:                                          ; preds = %entry
+  store i32 (i32)* @foo, i32 (i32)** %fp, align 8
+  br label %if.end
+
+if.else:                                          ; preds = %entry
+  store i32 (i32)* @bar, i32 (i32)** %fp, align 8
+  br label %if.end
+
+if.end:                                           ; preds = %if.else, %if.then
+  %1 = load i32 (i32)*, i32 (i32)** %fp, align 8
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  %call2 = call spir_func i64 @_Z13get_global_idj(i32 0) #4
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %call2
+  %3 = load i32, i32 addrspace(1)* %arrayidx, align 4
+  %call3 = call spir_func i32 @helper(i32 (i32)* %1, i32 %3) #3
+  %4 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  %call4 = call spir_func i64 @_Z13get_global_idj(i32 0) #4
+  %arrayidx5 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %call4
+  store i32 %call3, i32 addrspace(1)* %arrayidx5, align 4
+  ret void
+}
+
+; Function Attrs: convergent nounwind readnone
+declare spir_func i64 @_Z13get_global_idj(i32) #2
+
+attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #2 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #3 = { convergent }
+attributes #4 = { convergent nounwind readnone }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 0}
+!2 = !{i32 1, i32 2}
+!3 = !{!"clang version 7.1.0 "}
+!4 = !{!"none", !"none"}
+!5 = !{!"int*", !"int"}
+!6 = !{!"", !""}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll
new file mode 100644
index 0000000000..2ebbb429b3
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer-dedicated-as.ll
@@ -0,0 +1,107 @@
+; This test checks how a function pointer in a dedicated addr space would be
+; translated with and without -spirv-emit-function-ptr-addr-space option.
+; Expected behaviour:
+; No option is passed to the forward translation stage - no CodeSectionINTEL storage class in SPIR-V
+; The option is passed to the forward translation stage - CodeSectionINTEL storage class is generated
+; No option is passed to the reverse translation stage - function pointers are in private address space
+; The option is passed to the reverse translation stage - function pointers are in addrspace(9)
+;
+; Overall IR generation is tested elsewhere, here checks are very simple
+
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-AS
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spv
+; RUN: llvm-spirv -r %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-NO-AS
+
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-AS
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -spirv-emit-function-ptr-addr-space -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-AS
+
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-NO-AS
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-NO-AS
+
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-NO-AS
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM-AS
+
+; CHECK-SPIRV-AS-DAG: TypePointer [[#PtrCodeTy:]] 5605 [[#]]
+; CHECK-SPIRV-AS-DAG: TypePointer [[#PtrPrivTy:]] 7 [[#PtrCodeTy]]
+; CHECK-SPIRV-AS-DAG: ConstantFunctionPointerINTEL [[#PtrCodeTy]] [[#FunPtr:]]
+; CHECK-SPIRV-AS: Variable [[#PtrPrivTy]] [[#Var:]] 7
+; CHECK-SPIRV-AS: Store [[#Var]] [[#FunPtr]]
+; CHECK-SPIRV-AS: Load [[#PtrCodeTy]] [[#Load:]] [[#Var]]
+; CHECK-SPIRV-AS: FunctionPointerCallINTEL [[#]] [[#]] [[#Load]] [[#]]
+
+; CHECK-SPIRV-NO-AS-NOT: TypePointer [[#]] 5605 [[#]]
+
+; CHECK-LLVM-AS:  define spir_func i32 @foo(i32 %{{.*}}) addrspace(9)
+
+; CHECK-LLVM-NO-AS-NOT: addrspace(9)
+
+; ModuleID = 'function-pointer-dedicated-as.bc'
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-G1"
+target triple = "spir64-unknown-unknown"
+
+; Function Attrs: noinline nounwind
+define spir_func i32 @foo(i32 %arg) addrspace(9) #0 {
+entry:
+  %arg.addr = alloca i32, align 4
+  store i32 %arg, i32* %arg.addr, align 4
+  %0 = load i32, i32* %arg.addr, align 4
+  %add = add nsw i32 %0, 10
+  ret i32 %add
+}
+
+; Function Attrs: noinline nounwind
+define spir_kernel void @test(i32 addrspace(1)* %data, i32 %input) #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !6 !kernel_arg_type !7 !kernel_arg_type_qual !8 !kernel_arg_base_type !7 !spirv.ParameterDecorations !9 {
+entry:
+  %data.addr = alloca i32 addrspace(1)*, align 8
+  %input.addr = alloca i32, align 4
+  %fp = alloca i32 (i32) addrspace(9)*, align 8
+  store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8
+  store i32 %input, i32* %input.addr, align 4
+  store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %fp, align 8
+  %0 = load i32 (i32)  addrspace(9)*, i32 (i32)  addrspace(9)** %fp, align 8
+  %1 = load i32, i32* %input.addr, align 4
+  %call = call spir_func addrspace(9) i32 %0(i32 %1)
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  store i32 %call, i32 addrspace(1)* %2, align 4
+  ret void
+}
+
+attributes #0 = { noinline nounwind }
+
+!spirv.MemoryModel = !{!0}
+!spirv.Source = !{!1}
+!opencl.spir.version = !{!2}
+!opencl.ocl.version = !{!3}
+!opencl.used.extensions = !{!4}
+!opencl.used.optional.core.features = !{!4}
+!spirv.Generator = !{!5}
+
+!0 = !{i32 2, i32 2}
+!1 = !{i32 3, i32 100000}
+!2 = !{i32 1, i32 2}
+!3 = !{i32 1, i32 0}
+!4 = !{}
+!5 = !{i16 6, i16 14}
+!6 = !{!"none", !"none"}
+!7 = !{!"int*", !"int"}
+!8 = !{!"", !""}
+!9 = !{!4, !4}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll
new file mode 100644
index 0000000000..d43497b59f
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/function-pointer.ll
@@ -0,0 +1,92 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+;
+; Generated from:
+; int foo(int arg) {
+;   return arg + 10;
+; }
+;
+; void __kernel test(__global int *data, int input) {
+;   int (__constant *fp)(int) = &foo;
+;
+;   *data = fp(input);
+; }
+;
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: TypeInt [[TYPE_INT_ID:[0-9]+]]
+; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT_ID]] [[TYPE_INT_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_PTR_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_PTR_ALLOCA_ID:[0-9]+]] 7 [[FOO_PTR_ID]]
+; CHECK-SPIRV: ConstantFunctionPointerINTEL [[FOO_PTR_ID]] [[FOO_PTR:[0-9]+]] [[FOO_ID:[0-9]+]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] {{[0-9]+}} [[FOO_TYPE_ID]]
+; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]]
+; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_ID]] [[FOO_PTR_ALLOCA:[0-9]+]]
+; CHECK-SPIRV: Store [[FOO_PTR_ALLOCA]] [[FOO_PTR]]
+; CHECK-SPIRV: Load [[FOO_PTR_ID]] [[LOADED_FOO_PTR:[0-9]+]] [[FOO_PTR_ALLOCA]]
+; CHECK-SPIRV: FunctionPointerCallINTEL 2 {{[0-9]+}} [[LOADED_FOO_PTR]]
+;
+; CHECK-LLVM: define spir_kernel void @test
+; CHECK-LLVM: %fp = alloca i32 (i32) addrspace(9)*
+; CHECK-LLVM: store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %fp
+; CHECK-LLVM: %0 = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %fp
+; CHECK-LLVM: %call = call spir_func addrspace(9) i32 %0(i32 %1)
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @foo(i32 %arg) #0 {
+entry:
+  %arg.addr = alloca i32, align 4
+  store i32 %arg, i32* %arg.addr, align 4
+  %0 = load i32, i32* %arg.addr, align 4
+  %add = add nsw i32 %0, 10
+  ret i32 %add
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_kernel void @test(i32 addrspace(1)* %data, i32 %input) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
+entry:
+  %data.addr = alloca i32 addrspace(1)*, align 8
+  %input.addr = alloca i32, align 4
+  %fp = alloca i32 (i32)*, align 8
+  store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8
+  store i32 %input, i32* %input.addr, align 4
+  store i32 (i32)* @foo, i32 (i32)** %fp, align 8
+  %0 = load i32 (i32)*, i32 (i32)** %fp, align 8
+  %1 = load i32, i32* %input.addr, align 4
+  %call = call spir_func i32 %0(i32 %1) #2
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  store i32 %call, i32 addrspace(1)* %2, align 4
+  ret void
+}
+
+attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #2 = { convergent }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!2}
+!opencl.used.extensions = !{!3}
+!opencl.used.optional.core.features = !{!3}
+!opencl.compiler.options = !{!3}
+!llvm.ident = !{!4}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 0}
+!2 = !{i32 1, i32 2}
+!3 = !{}
+!4 = !{!"clang version 7.0.0 "}
+!5 = !{!"none", !"none"}
+!6 = !{!"int*", !"int"}
+!7 = !{!"", !""}
+
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll
new file mode 100644
index 0000000000..31c36f98eb
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/gv-func-ptr.ll
@@ -0,0 +1,40 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+
+; ModuleID = 't.bc'
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024"
+target triple = "spir64-unknown-unknown"
+
+%structtype.3 = type { [1 x i8 addrspace(4)*] }
+
+; CHECK-LLVM: @A = addrspace(1) constant %structtype.3 { [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8 addrspace(9)* bitcast (void () addrspace(9)* @foo to i8 addrspace(9)*) to i8 addrspace(4)*)] }, align 8
+
+@A = linkonce_odr addrspace(1) constant %structtype.3 { [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8* bitcast (void ()* @foo to i8*) to i8 addrspace(4)*)] }, align 8
+
+; Function Attrs: nounwind
+define linkonce_odr spir_func void @foo() #0 {
+entry:
+; CHECK-LLVM: %0 = getelementptr inbounds %structtype.3, %structtype.3 addrspace(1)* @A, i64 0, i32 0, i64 2
+  %0 = getelementptr inbounds %structtype.3, %structtype.3 addrspace(1)* @A, i64 0, i32 0, i64 2
+  ret void
+}
+
+attributes #0 = { nounwind }
+
+!spirv.MemoryModel = !{!0}
+!spirv.Source = !{!1}
+!opencl.spir.version = !{!2}
+!opencl.ocl.version = !{!2}
+!opencl.used.extensions = !{!3}
+!opencl.used.optional.core.features = !{!4}
+!spirv.Generator = !{!5}
+
+!0 = !{i32 2, i32 2}
+!1 = !{i32 4, i32 200000}
+!2 = !{i32 2, i32 0}
+!3 = !{!"cl_khr_int64_extended_atomics", !"cl_khr_subgroups"}
+!4 = !{!"cl_doubles"}
+!5 = !{i16 6, i16 14}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll
new file mode 100644
index 0000000000..b98fb602f4
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/non-uniform-function-pointer.ll
@@ -0,0 +1,139 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+;
+; Generated from:
+; int foo(int v) {
+;   return v + 1;
+; }
+;
+; int bar(int v) {
+;   return v + 2;
+; }
+;
+; __kernel void test(__global int *data, int control) {
+;   int (*fp)(int) = 0;
+;
+;   if (get_global_id(0) % control == 0)
+;     fp = &foo;
+;   else
+;     fp = &bar;
+;
+;   data[get_global_id(0)] = fp(data[get_global_id(0)]);
+; }
+;
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+;
+; CHECK-SPIRV: Name [[KERNEL_ID:[0-9]+]] "test"
+; CHECK-SPIRV: TypeInt [[TYPE_INT32_ID:[0-9+]]] 32
+; CHECK-SPIRV: TypeFunction [[FOO_TYPE_ID:[0-9]+]] [[TYPE_INT32_ID]] [[TYPE_INT32_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_PTR_TYPE_ID:[0-9]+]] {{[0-9]+}} [[FOO_TYPE_ID]]
+; CHECK-SPIRV: TypePointer [[FOO_PTR_ALLOCA_TYPE_ID:[0-9]+]] 7 [[FOO_PTR_TYPE_ID]]
+; CHECK-SPIRV: ConstantFunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[FOO_PTR_ID:[0-9]+]] [[FOO_ID:[0-9]+]]
+; CHECK-SPIRV: ConstantFunctionPointerINTEL [[FOO_PTR_TYPE_ID]] [[BAR_PTR_ID:[0-9]+]] [[BAR_ID:[0-9]+]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]] {{[0-9]+}} [[FOO_TYPE_ID]]
+; CHECK-SPIRV: Function {{[0-9]+}} [[BAR_ID]] {{[0-9]+}} [[FOO_TYPE_ID]]
+;
+; CHECK-SPIRV: Function {{[0-9]+}} [[KERNEL_ID]]
+; CHECK-SPIRV: Variable [[FOO_PTR_ALLOCA_TYPE_ID]] [[FOO_PTR_ALLOCA_ID:[0-9]+]]
+; CHECK-SPIRV: Store [[FOO_PTR_ALLOCA_ID]] [[FOO_PTR_ID]]
+; CHECK-SPIRV: Store [[FOO_PTR_ALLOCA_ID]] [[BAR_PTR_ID]]
+; CHECK-SPIRV: Load [[FOO_PTR_TYPE_ID]] [[LOADED_FOO_PTR:[0-9]+]] [[FOO_PTR_ALLOCA_ID]]
+; CHECK-SPIRV: FunctionPointerCallINTEL {{[0-9]+}} {{[0-9]+}} [[LOADED_FOO_PTR]]
+;
+; CHECK-LLVM: define spir_kernel void @test
+; CHECK-LLVM: %fp = alloca i32 (i32) addrspace(9)*
+; CHECK-LLVM: store i32 (i32) addrspace(9)* @foo, i32 (i32) addrspace(9)** %fp
+; CHECK-LLVM: store i32 (i32) addrspace(9)* @bar, i32 (i32) addrspace(9)** %fp
+; CHECK-LLVM: %[[FP:.*]] = load i32 (i32) addrspace(9)*, i32 (i32) addrspace(9)** %fp
+; CHECK-LLVM: call spir_func addrspace(9) i32 %[[FP]](i32 %{{.*}})
+
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @foo(i32 %v) #0 {
+entry:
+  %v.addr = alloca i32, align 4
+  store i32 %v, i32* %v.addr, align 4
+  %0 = load i32, i32* %v.addr, align 4
+  %add = add nsw i32 %0, 1
+  ret i32 %add
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @bar(i32 %v) #0 {
+entry:
+  %v.addr = alloca i32, align 4
+  store i32 %v, i32* %v.addr, align 4
+  %0 = load i32, i32* %v.addr, align 4
+  %add = add nsw i32 %0, 2
+  ret i32 %add
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_kernel void @test(i32 addrspace(1)* %data, i32 %control) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 {
+entry:
+  %data.addr = alloca i32 addrspace(1)*, align 8
+  %control.addr = alloca i32, align 4
+  %fp = alloca i32 (i32)*, align 8
+  store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8
+  store i32 %control, i32* %control.addr, align 4
+  store i32 (i32)* null, i32 (i32)** %fp, align 8
+  %call = call spir_func i64 @_Z13get_global_idj(i32 0) #3
+  %0 = load i32, i32* %control.addr, align 4
+  %conv = sext i32 %0 to i64
+  %rem = urem i64 %call, %conv
+  %cmp = icmp eq i64 %rem, 0
+  br i1 %cmp, label %if.then, label %if.else
+
+if.then:                                          ; preds = %entry
+  store i32 (i32)* @foo, i32 (i32)** %fp, align 8
+  br label %if.end
+
+if.else:                                          ; preds = %entry
+  store i32 (i32)* @bar, i32 (i32)** %fp, align 8
+  br label %if.end
+
+if.end:                                           ; preds = %if.else, %if.then
+  %1 = load i32 (i32)*, i32 (i32)** %fp, align 8
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  %call2 = call spir_func i64 @_Z13get_global_idj(i32 0) #3
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %call2
+  %3 = load i32, i32 addrspace(1)* %arrayidx, align 4
+  %call3 = call spir_func i32 %1(i32 %3) #4
+  %4 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  %call4 = call spir_func i64 @_Z13get_global_idj(i32 0) #3
+  %arrayidx5 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %call4
+  store i32 %call3, i32 addrspace(1)* %arrayidx5, align 4
+  ret void
+}
+
+; Function Attrs: convergent nounwind readnone
+declare spir_func i64 @_Z13get_global_idj(i32) #2
+
+attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #2 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #3 = { convergent nounwind readnone }
+attributes #4 = { convergent }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 0}
+!2 = !{i32 1, i32 2}
+!3 = !{!"clang version 7.1.0 "}
+!4 = !{!"none", !"none"}
+!5 = !{!"int*", !"int"}
+!6 = !{!"", !""}
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll
new file mode 100644
index 0000000000..2d7d0dff76
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/referenced-indirectly.ll
@@ -0,0 +1,82 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc -spirv-text --spirv-ext=+SPV_INTEL_function_pointers -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+;
+; Generated from:
+; __attribute__((referenced_indirectly))
+; int foo(int arg) {
+;   return arg + 10;
+; }
+;
+; void __kernel test(__global int *data, int input) {
+;   int (__constant *fp)(int) = &foo;
+;
+;   *data = fp(input);
+; }
+;
+; CHECK-SPIRV: Capability FunctionPointersINTEL
+; CHECK-SPIRV: Capability IndirectReferencesINTEL
+; CHECK-SPIRV: Extension "SPV_INTEL_function_pointers"
+;
+; CHECK-SPIRV: Name [[FOO_ID:[0-9]+]] "foo"
+; CHECK-SPIRV: Decorate [[FOO_ID]] ReferencedIndirectlyINTEL
+; CHECK-SPIRV: Function {{[0-9]+}} [[FOO_ID]]
+;
+; CHECK-LLVM: define spir_func i32 @foo(i32 %arg) addrspace(9) #[[ATTRS:[0-9]+]]
+; CHECK-LLVM: attributes #[[ATTRS]] = {{.*}} "referenced-indirectly"
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_func i32 @foo(i32 %arg) #0 {
+entry:
+  %arg.addr = alloca i32, align 4
+  store i32 %arg, i32* %arg.addr, align 4
+  %0 = load i32, i32* %arg.addr, align 4
+  %add = add nsw i32 %0, 10
+  ret i32 %add
+}
+
+; Function Attrs: convergent noinline nounwind optnone
+define spir_kernel void @test(i32 addrspace(1)* %data, i32 %input) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
+entry:
+  %data.addr = alloca i32 addrspace(1)*, align 8
+  %input.addr = alloca i32, align 4
+  %fp = alloca i32 (i32)*, align 8
+  store i32 addrspace(1)* %data, i32 addrspace(1)** %data.addr, align 8
+  store i32 %input, i32* %input.addr, align 4
+  store i32 (i32)* @foo, i32 (i32)** %fp, align 8
+  %0 = load i32 (i32)*, i32 (i32)** %fp, align 8
+  %1 = load i32, i32* %input.addr, align 4
+  %call = call spir_func i32 %0(i32 %1) #2
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %data.addr, align 8
+  store i32 %call, i32 addrspace(1)* %2, align 4
+  ret void
+}
+
+attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" "referenced-indirectly" }
+attributes #1 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #2 = { convergent }
+
+!llvm.module.flags = !{!0}
+!opencl.ocl.version = !{!1}
+!opencl.spir.version = !{!2}
+!opencl.used.extensions = !{!3}
+!opencl.used.optional.core.features = !{!3}
+!opencl.compiler.options = !{!3}
+!llvm.ident = !{!4}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 0}
+!2 = !{i32 1, i32 2}
+!3 = !{}
+!4 = !{!"clang version 7.0.0 "}
+!5 = !{!"none", !"none"}
+!6 = !{!"int*", !"int"}
+!7 = !{!"", !""}
+
diff --git a/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll
new file mode 100644
index 0000000000..020c93a996
--- /dev/null
+++ b/test/transcoding/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll
@@ -0,0 +1,147 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_function_pointers -o %t.spv
+; RUN: llvm-spirv %t.spv -to-text -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+; RUN: llvm-spirv -r -spirv-emit-function-ptr-addr-space %t.spv -o %t.r.bc
+; RUN: llvm-dis %t.r.bc -o %t.r.ll
+; RUN: FileCheck < %t.r.ll %s --check-prefix=CHECK-LLVM
+
+; CHECK-SPIRV: Name [[#KERNEL_ID:]] "_ZTS6kernel"
+; CHECK-SPIRV-DAG: Name [[#BAR:]] "_Z3barii"
+; CHECK-SPIRV-DAG: Name [[#BAZ:]] "_Z3bazii"
+; CHECK-SPIRV: TypeInt [[#INT32:]] 32
+; CHECK-SPIRV: TypeFunction [[#FUNC_TYPE:]] [[#INT32]] [[#INT32]]
+; CHECK-SPIRV: TypePointer [[#FUNC_PTR_TYPE:]] [[#]] [[#FUNC_TYPE]]
+; CHECK-SPIRV: TypePointer [[#FUNC_PTR_ALLOCA_TYPE:]] [[#]] [[#FUNC_PTR_TYPE]]
+; CHECK-SPIRV-DAG: ConstantFunctionPointerINTEL [[#FUNC_PTR_TYPE]] [[#BARPTR:]] [[#BAR]]
+; CHECK-SPIRV-DAG: ConstantFunctionPointerINTEL [[#FUNC_PTR_TYPE]] [[#BAZPTR:]] [[#BAZ]]
+; CHECK-SPIRV: Function [[#]] [[#KERNEL_ID]]
+; CHECK-SPIRV: Variable [[#FUNC_PTR_ALLOCA_TYPE]] [[#FPTR:]]
+; CHECK-SPIRV: Select [[#FUNC_PTR_TYPE]] [[#SELECT:]] [[#]] [[#BARPTR]] [[#BAZPTR]]
+; CHECK-SPIRV: Store [[#FPTR]] [[#SELECT]]
+; CHECK-SPIRV: Load [[#FUNC_PTR_TYPE]] [[#LOAD:]] [[#FPTR]]
+; CHECK-SPIRV: FunctionPointerCallINTEL [[#]] [[#]] [[#LOAD]]
+
+; CHECK-LLVM: define spir_kernel void @_ZTS6kernel
+; CHECK-LLVM: %[[FPTR_ALLOCA:.*]] = alloca i32 (i32, i32) addrspace(9)*
+; CHECK-LLVM: %[[SELECT:.*]] = select i1 %{{.*}}, i32 (i32, i32) addrspace(9)* @_Z3barii, i32 (i32, i32) addrspace(9)* @_Z3bazii
+; CHECK-LLVM: store i32 (i32, i32) addrspace(9)* %[[SELECT]], i32 (i32, i32) addrspace(9)** %[[FPTR_ALLOCA]]
+; CHECK-LLVM: %[[FPTR:.*]] = load i32 (i32, i32) addrspace(9)*, i32 (i32, i32) addrspace(9)** %[[FPTR_ALLOCA]]
+; CHECK-LLVM: call spir_func addrspace(9) i32 %[[FPTR]](
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
+%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
+
+$_ZTS6kernel = comdat any
+
+@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+
+; Function Attrs: norecurse
+define weak_odr dso_local spir_kernel void @_ZTS6kernel(i32 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
+entry:
+  %fptr.alloca = alloca i32 (i32, i32)*, align 8
+  %ref.tmp.i = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8
+  %agg.tmp2.i = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8
+  %agg.tmp3.i = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8
+  %agg.tmp6 = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8
+  %0 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp2.i to i8*
+  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %0)
+  %1 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp3.i to i8*
+  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %1)
+  %2 = addrspacecast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp2.i to %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)*
+  %ptrint4.i = ptrtoint %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %2 to i64
+  %maskedptr5.i = and i64 %ptrint4.i, 7
+  %maskcond6.i = icmp eq i64 %maskedptr5.i, 0
+  %3 = addrspacecast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp3.i to %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)*
+  %ptrint.i = ptrtoint %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %3 to i64
+  %maskedptr.i = and i64 %ptrint.i, 7
+  %maskcond.i = icmp eq i64 %maskedptr.i, 0
+  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %0)
+  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %1)
+  %4 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
+  %5 = load i64, i64* %4, align 8
+  %add.ptr.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_, i64 %5
+  %6 = addrspacecast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp6 to %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)*
+  %ptrint = ptrtoint %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %6 to i64
+  %maskedptr = and i64 %ptrint, 7
+  %maskcond = icmp eq i64 %maskedptr, 0
+  %7 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !8
+  %8 = extractelement <3 x i64> %7, i64 0
+  %arrayinit.begin.i.i.i.i.i = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %6, i64 0, i32 0, i32 0, i64 0
+  store i64 %8, i64 addrspace(4)* %arrayinit.begin.i.i.i.i.i, align 8, !tbaa !15, !alias.scope !8
+  %9 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %ref.tmp.i to i8*
+  call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %9) #4
+  %10 = addrspacecast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %ref.tmp.i to %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)*
+  %ptrint.i2 = ptrtoint %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %10 to i64
+  %maskedptr.i3 = and i64 %ptrint.i2, 7
+  %maskcond.i4 = icmp eq i64 %maskedptr.i3, 0
+  %rem.i.i = and i64 %8, 1
+  %cmp.i.i = icmp eq i64 %rem.i.i, 0
+  call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %9) #4
+  %_Z3barii._Z3bazii.i = select i1 %cmp.i.i, i32 (i32, i32)* @_Z3barii, i32 (i32, i32)* @_Z3bazii
+  store i32 (i32, i32)* %_Z3barii._Z3bazii.i, i32 (i32, i32)** %fptr.alloca, align 8
+  %fptr = load i32 (i32, i32)*, i32 (i32, i32)** %fptr.alloca, align 8
+  %call4.i = call spir_func i32 %fptr(i32 10, i32 10), !callees !19
+  %arrayidx.i3.i = getelementptr inbounds i32, i32 addrspace(1)* %add.ptr.i, i64 %8
+  %arrayidx.ascast.i.i = addrspacecast i32 addrspace(1)* %arrayidx.i3.i to i32 addrspace(4)*
+  store i32 %call4.i, i32 addrspace(4)* %arrayidx.ascast.i.i, align 4, !tbaa !20
+  ret void
+}
+
+; Function Attrs: argmemonly nounwind willreturn
+declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
+
+; Function Attrs: argmemonly nounwind willreturn
+declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
+
+; Function Attrs: norecurse nounwind readnone
+define dso_local spir_func i32 @_Z3barii(i32 %a, i32 %b) local_unnamed_addr #2 {
+entry:
+  %add = add nsw i32 %b, %a
+  ret i32 %add
+}
+
+; Function Attrs: norecurse nounwind readnone
+define dso_local spir_func i32 @_Z3bazii(i32 %a, i32 %b) local_unnamed_addr #2 {
+entry:
+  %sub = sub nsw i32 %a, %b
+  ret i32 %sub
+}
+
+attributes #0 = { norecurse "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "sycl-module-id"="f.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #1 = { argmemonly nounwind willreturn }
+attributes #2 = { norecurse nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
+attributes #3 = { nounwind willreturn }
+attributes #4 = { nounwind }
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+!spirv.Source = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 11.0.0 "}
+!4 = !{i32 1, i32 0, i32 0, i32 0}
+!5 = !{!"none", !"none", !"none", !"none"}
+!6 = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"}
+!7 = !{!"", !"", !"", !""}
+!8 = !{!9, !11, !13}
+!9 = distinct !{!9, !10, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"}
+!10 = distinct !{!10, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"}
+!11 = distinct !{!11, !12, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"}
+!12 = distinct !{!12, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"}
+!13 = distinct !{!13, !14, !"_ZN2cl4sycl6detail7Builder5getIdILi1EEEKNS0_2idIXT_EEEv: %agg.result"}
+!14 = distinct !{!14, !"_ZN2cl4sycl6detail7Builder5getIdILi1EEEKNS0_2idIXT_EEEv"}
+!15 = !{!16, !16, i64 0}
+!16 = !{!"long", !17, i64 0}
+!17 = !{!"omnipotent char", !18, i64 0}
+!18 = !{!"Simple C++ TBAA"}
+!19 = !{i32 (i32, i32)* @_Z3barii, i32 (i32, i32)* @_Z3bazii}
+!20 = !{!21, !21, i64 0}
+!21 = !{!"int", !17, i64 0}
diff --git a/tools/llvm-spirv/llvm-spirv.cpp b/tools/llvm-spirv/llvm-spirv.cpp
index 2a09c2f132..fadd3b9cec 100644
--- a/tools/llvm-spirv/llvm-spirv.cpp
+++ b/tools/llvm-spirv/llvm-spirv.cpp
@@ -161,6 +161,10 @@ static cl::opt<bool>
                                 "for the translation from SPIR-V."),
                        cl::Hidden);
 
+static cl::opt<bool> SPIRVEmitFunctionPtrAddrSpace(
+    "spirv-emit-function-ptr-addr-space", cl::init(false),
+    cl::desc("Emit and consume CodeSectionINTEL for function pointers"));
+
 using SPIRV::ExtensionID;
 
 #ifdef _SPIRV_SUPPORT_TEXT_FMT
@@ -767,6 +771,9 @@ int main(int Ac, char **Av) {
   if (PreserveOCLKernelArgTypeMetadataThroughString.getNumOccurrences() != 0)
     Opts.setPreserveOCLKernelArgTypeMetadataThroughString(true);
 
+  if (SPIRVEmitFunctionPtrAddrSpace.getNumOccurrences() != 0)
+    Opts.setEmitFunctionPtrAddrSpace(true);
+
 #ifdef _SPIRV_SUPPORT_TEXT_FMT
   if (ToText && (ToBinary || IsReverse || IsRegularization)) {
     errs() << "Cannot use -to-text with -to-binary, -r, -s\n";