TGSI: Fix isKernelFunction

Signed-off-by: Hans de Goede <hdegoede@redhat.com>

2 files changed, 16 insertions, 20 deletions

diff --git a/lib/Target/TGSI/TGSI.h b/lib/Target/TGSI/TGSI.h
index c906c8f98ad..478f6c69f82 100644
--- a/lib/Target/TGSI/TGSI.h
+++ b/lib/Target/TGSI/TGSI.h
@@ -42,30 +42,26 @@ namespace llvm {
       };
    }
 
-#if 0
    namespace {
-      inline bool isKernelFunction(const Function *f) {
-         NamedMDNode *md = f->getParent()->getNamedMetadata("opencl.kernels");
-
-         if (md) {
-            printf("isKernelFunction found md with %d operands\n");
-            for (unsigned i = 0; i < md->getNumOperands(); ++i) {
-	       auto *metaData = dyn_cast<ConstantAsMetadata>(md->getOperand(i)->getOperand(0));
-
-	       printf("isKernelFunction md %d %p\n", i, metaData);
-	       assert(metaData);
-
-	       Value *v = metaData->getValue();
-               assert(v && isa<Function>(v));
+      inline const MDNode *getKernelMetadata(const Function *f) {
+         const Module *m = f->getParent();
+         assert(m);
+         NamedMDNode *nmd = m->getNamedMetadata("opencl.kernels");
+         assert(nmd);
 
-               if (f == static_cast<Function *>(v))
-                  return true;
-            }
+         for (unsigned i = 0; i < nmd->getNumOperands(); ++i) {
+            const MDNode *md = nmd->getOperand(i);
+            const Function *md_f = mdconst::dyn_extract<llvm::Function>(
+               md->getOperand(0));
+            if (md_f == f)
+               return md;
          }
+         return NULL;
+      }
 
-         return false;
+      inline bool isKernelFunction(const Function *f) {
+         return getKernelMetadata(f) != NULL;
       }
    }
-#endif
 }
 #endif
diff --git a/lib/Target/TGSI/TGSITargetTransformInfo.cpp b/lib/Target/TGSI/TGSITargetTransformInfo.cpp
index 7c4d0d6883f..6811425962a 100644
--- a/lib/Target/TGSI/TGSITargetTransformInfo.cpp
+++ b/lib/Target/TGSI/TGSITargetTransformInfo.cpp
@@ -24,7 +24,7 @@ bool TGSITTIImpl::isSourceOfDivergence(const Value *V) {
   // Without inter-procedural analysis, we conservatively assume that arguments
   // to __device__ functions are divergent.
   if (const Argument *Arg = dyn_cast<Argument>(V))
-    return true; // FIXME: !isKernelFunction(*Arg->getParent());
+    return !isKernelFunction(Arg->getParent());
 
   if (const Instruction *I = dyn_cast<Instruction>(V)) {
     // Without pointer analysis, we conservatively assume values loaded from