blob: 6811425962ac8e16cdc63c0586f62b6ddab011e8 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
|
//===---- TGSITargetTransformInfo.cpp - TGSI specific TTI -----------------===//
//
// The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
#include "TGSITargetTransformInfo.h"
#include "llvm/Analysis/LoopInfo.h"
#include "llvm/Analysis/TargetTransformInfo.h"
#include "llvm/Analysis/ValueTracking.h"
#include "llvm/CodeGen/BasicTTIImpl.h"
#include "llvm/Support/Debug.h"
#include "llvm/Target/CostTable.h"
#include "llvm/Target/TargetLowering.h"
using namespace llvm;
#define DEBUG_TYPE "TGSItti"
// Based in the NVPTX isSourceOfDivergence() implementation
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 !isKernelFunction(Arg->getParent());
if (const Instruction *I = dyn_cast<Instruction>(V)) {
// Without pointer analysis, we conservatively assume values loaded from
// generic or local address space are divergent.
if (const LoadInst *LI = dyn_cast<LoadInst>(I)) {
unsigned AS = LI->getPointerAddressSpace();
return AS == tgsi::PRIVATE || AS == tgsi::LOCAL;
}
// Atomic instructions may cause divergence. Atomic instructions are
// executed sequentially across all threads in a warp. Therefore, an earlier
// executed thread may see different memory inputs than a later executed
// thread. For example, suppose *a = 0 initially.
//
// atom.global.add.s32 d, [a], 1
//
// returns 0 for the first thread that enters the critical region, and 1 for
// the second thread.
if (I->isAtomic())
return true;
if (const IntrinsicInst *II = dyn_cast<IntrinsicInst>(I)) {
switch (II->getIntrinsicID()) {
case Intrinsic::tgsi_read_blockid_x:
case Intrinsic::tgsi_read_blockid_y:
case Intrinsic::tgsi_read_blockid_z:
case Intrinsic::tgsi_read_threadid_x:
case Intrinsic::tgsi_read_threadid_y:
case Intrinsic::tgsi_read_threadid_z:
return true;
default:
return false;
}
}
// Conservatively consider the return value of function calls as divergent.
// We could analyze callees with bodies more precisely using
// inter-procedural analysis.
if (isa<CallInst>(V) || isa<InvokeInst>(V))
return true;
}
return false;
}
|
