search for: swpenim

...real" ISAs, these structures do not contain any special properties about > registers or instructions, like cost or scheduling information. Are you > trying to figure out the total number of PTX registers that will be emitted? > > > On Wed, Aug 7, 2013 at 7:06 AM, Anthony Yu <swpenim at gmail.com> wrote: > >> OK. I know what you mean......... >> >> Simply speaking, I want to do some optimizations for PTX, and the >> information I need is similar to a register allocator. I know PTX is >> virtual ISA, but I will use PTX as the input of the simu...

...node's value would be the shared address space pointer of the allocation, and the assembly printer would turn that into a ".shared .bX ..." variable. Would that solve your issue? Or do you need to change other parts of the IR as well? On Thu, Aug 8, 2013 at 7:29 AM, Anthony Yu <swpenim at gmail.com> wrote: > Yes, total number of PTX registers that will be emitted is exactly what I > need. It's hard to figure out this in LLVM IR level. > > > 2013/8/7 Justin Holewinski <justin.holewinski at gmail.com> > >> Is there any way you could approximate...

...nctionpass-class > ). > > If you explain what you're trying to accomplish, I can try to help figure > out a good approach here. There very well may be limitations to the > current infrastructure that need to be fixed. > > > On Tue, Aug 6, 2013 at 1:25 AM, Anthony Yu <swpenim at gmail.com> wrote: > >> I want to create share memory in my MachineFunctionPass, and insert >> load/store instruction for it. The way to create share memory is to add >> global variables which are in share memory address space (not sure if it is >> the only way). The...

...ends that target "real" ISAs, these structures do not contain any special properties about registers or instructions, like cost or scheduling information. Are you trying to figure out the total number of PTX registers that will be emitted? On Wed, Aug 7, 2013 at 7:06 AM, Anthony Yu <swpenim at gmail.com> wrote: > OK. I know what you mean......... > > Simply speaking, I want to do some optimizations for PTX, and the > information I need is similar to a register allocator. I know PTX is > virtual ISA, but I will use PTX as the input of the simulator, gpgpu-sim, > s...

...n > will likely propagate down through code generation, but at that point what > is the purpose of using a MachineFunctionPass? You won't have any analysis > or instruction information available until runOnMachineFunction. > > > On Mon, Aug 5, 2013 at 12:00 PM, Anthony Yu <swpenim at gmail.com> wrote: > >> Micah, >> >> Thanks for your help. I will study on that code. >> >> >> Justin, >> >> Sorry for my misleading word. Local memory in OpenCL is the same as share >> memory in CUDA. What I mean is share memory, so Mac...

...s/WritingAnLLVMPass.html#the-machinefunctionpass-class). If you explain what you're trying to accomplish, I can try to help figure out a good approach here. There very well may be limitations to the current infrastructure that need to be fixed. On Tue, Aug 6, 2013 at 1:25 AM, Anthony Yu <swpenim at gmail.com> wrote: > I want to create share memory in my MachineFunctionPass, and insert > load/store instruction for it. The way to create share memory is to add > global variables which are in share memory address space (not sure if it is > the only way). Therefore, I should add...

...alue of these methods is an integer that represents a > FrameIndex. You can treat this as a pointer to your allocated object. You > will also need to emit the proper MachineInstr-level loads and stores to > access the object. > > > On Mon, Aug 5, 2013 at 11:00 AM, Anthony Yu <swpenim at gmail.com> wrote: > >> Micah, >> >> As you expected, I am trying to create local memory but in the NVPTX >> backend. It's really not convenient that I can't create local memory in >> runOnMachineFunction. >> Hmm.... >> Since I should do it...

...unctionPass::doInitialization will likely propagate down through code generation, but at that point what is the purpose of using a MachineFunctionPass? You won't have any analysis or instruction information available until runOnMachineFunction. On Mon, Aug 5, 2013 at 12:00 PM, Anthony Yu <swpenim at gmail.com> wrote: > Micah, > > Thanks for your help. I will study on that code. > > > Justin, > > Sorry for my misleading word. Local memory in OpenCL is the same as share > memory in CUDA. What I mean is share memory, so MachineFrameInfo is not > suitable to me...

Micah, As you expected, I am trying to create local memory but in the NVPTX backend. It's really not convenient that I can't create local memory in runOnMachineFunction. Hmm.... Since I should do it at doInitialization stage, I also need to do some tricks in global variable and AsmPrinter to resize it. Did you use the similar way? Antony 2013/8/5 Micah Villmow <micah.villmow at

...emory in PTX). The return value of these methods is an integer that represents a FrameIndex. You can treat this as a pointer to your allocated object. You will also need to emit the proper MachineInstr-level loads and stores to access the object. On Mon, Aug 5, 2013 at 11:00 AM, Anthony Yu <swpenim at gmail.com> wrote: > Micah, > > As you expected, I am trying to create local memory but in the NVPTX > backend. It's really not convenient that I can't create local memory in > runOnMachineFunction. > Hmm.... > Since I should do it at doInitialization stage, I als...

Address space 4 globals should be going into .const. Please open a bug report and provide a repro .ll file. On Fri, Jun 7, 2013 at 11:37 AM, Antony Yu <swpenim at gmail.com> wrote: > Hello, > > I work on compiling OpenCL kernel to PTX code by clang and NVPTX with > libclc. > I have a kernel that contains constant variable declared in file scope like > this: > > constant one_f = 1.0f; > __kernel void test( ...){ ... }...

Oops! No need of Call Graph! In fact, what I want to do is to find which function is the kernel function and which function is called by that kernel. Since OpenCL will make all functions called by kernels inline, I can use function attribute: Noinline to distinguish them. Sorry for bothering you. Antony Yu -- View this message in context:

You shouldn't rely on that, its an implementation detail. Instead, you can trace back to the original Function object and check for kernel metadata. See http://llvm.org/docs/NVPTXUsage.html#marking-functions-as-kernels On Tue, Jun 25, 2013 at 11:09 AM, Antony Yu <swpenim at gmail.com> wrote: > Oops! No need of Call Graph! > In fact, what I want to do is to find which function is the kernel function > and which function is called by that kernel. Since OpenCL will make all > functions called by kernels inline, I can use function attribute: Noinline &g...

Are you sure you are initializing your pass properly? Can you show a stripped down version of your pass? On Fri, Jun 21, 2013 at 7:27 AM, Anthony Yu <swpenim at gmail.com> wrote: > Hello, > > I want to write a modulePass in addPreEmitPass() for NVPTX, but I > encounter an assertion failed when executing clang. > > Here is my error message. > ==== > Pass 'NVPTX Assembly Printer' is not initialized. > Verify if there...

Hello, I work on compiling OpenCL kernel to PTX code by clang and NVPTX with libclc. I have a kernel that contains constant variable declared in file scope like this: constant one_f = 1.0f; __kernel void test( ...){ ... } Then it is compiled to llvm-ir: @one_f = addrspace(4) const float 1.000000e+00, align 4 define void test(...){ ... } Finally ptx: .visible .global .align 4 .f32

I try to use INITIALIZE_PASS instead of RegisterPass<> to register my pass, though I don't understand what's their difference and how it works because its documents doesn't exist. But it still doesn't work. Parts of my codes is as follows: in NVPTXTest.h namespace llvm { void initializeNVPTXTestPass(PassRegistry &r); class NVPTXTest : public ModulePass { public:

...ed to machine code, its too late to run LLVM IR passes (ModulePass, FunctionPass, CallGraphSCCPass, etc.). At that point, you need to run a Machine*Pass, e.g. MachineFunctionPass. If you need to run an IR level pass, you need to use addIRPasses(). On Mon, Jun 24, 2013 at 11:20 AM, Antony Yu <swpenim at gmail.com> wrote: > I try to use INITIALIZE_PASS instead of RegisterPass<> to register my pass, > though I don't understand what's their difference and how it works because > its documents doesn't exist. But it still doesn't work. > > Parts of my codes is...

I write my pass in a mix way of NVPTXAllocaHoisting, NVPTXSplitBBatBar and transforms/Hello. The following is part of the codes: in NVPTXTargetMachine.cpp bool NVPTXPassConfig::addPreEmitPass() { addPass(createTest()); return false; } in NVPTXTest.h namespace llvm{

Hello, I want to write a modulePass in addPreEmitPass() for NVPTX, but I encounter an assertion failed when executing clang. Here is my error message. ==== Pass 'NVPTX Assembly Printer' is not initialized. Verify if there is a pass dependency cycle. Required Passes: llc: /home/pyyu/local/llvm/lib/IR/PassManager.cpp:637: void llvm::PMTopLevelManager::schedulePass(llvm::Pass*): Assertion

So that's how it is !! Thanks! Hmm, in fact, I want to get Call Graph information in machine code stage, is there any way to get this? I try CallGraph &g= getAnalysis<CallGraph>, but it prints the same error and doesn't work as expected. Antony Yu Justin Holewinski-2 wrote > Sorry for the delay. Yeah, that error message is a bit confusing. What's > happening is