Let SPIRVWriter translate TypedPointerTypes directly. (#1629)

Author: jcranmer-intel

lib/SPIRV/SPIRVWriter.cpp

@@ -82,6 +82,7 @@
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/Operator.h"
+#include "llvm/IR/TypedPointerType.h"
 #include "llvm/Pass.h"
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/Support/Casting.h"
@@ -342,6 +343,10 @@ SPIRVType *LLVMToSPIRVBase::transType(Type *T) {
     return transPointerType(ET, AddrSpc);
   }
 
+  if (auto *TPT = dyn_cast<TypedPointerType>(T)) {
+    return transPointerType(TPT->getElementType(), TPT->getAddressSpace());
+  }
+
   if (auto *VecTy = dyn_cast<FixedVectorType>(T)) {
     if (VecTy->getElementType()->isPointerTy()) {
       // SPV_INTEL_masked_gather_scatter extension changes 2.16.1. Universal

test/event_no_group_cap.cl

@@ -2,7 +2,7 @@ __kernel void test_fn( const __global char *src)
 {
 	wait_group_events(0, NULL);
 }
-// RUN: %clang_cc1 -triple spir64 -x cl -cl-std=CL2.0 -fdeclare-opencl-builtins -finclude-default-header -O0 -emit-llvm-bc %s -o %t.bc -no-opaque-pointers
+// RUN: %clang_cc1 -triple spir64 -x cl -cl-std=CL2.0 -fdeclare-opencl-builtins -finclude-default-header -O0 -emit-llvm-bc %s -o %t.bc
 // RUN: llvm-spirv %t.bc -spirv-text -o %t.spt
 // RUN: FileCheck < %t.spt %s
 // RUN: llvm-spirv %t.bc -o %t.spv