search for: int_nvvm_read_ptx_sreg_tid_x

...<"__builtin_amdgcn_workgroup_id">; There is no new definition of any intrinsics within the target AMDGPU. I was working before with the NVPTX backend and that target has the special registers are associated with strings for the LLVM IR, e.g., PTX_READ_SREG_R32<"tid.x", int_nvvm_read_ptx_sreg_tid_x>; Maybe I'm missing something but does this mean at the moment it is not possible to write a function in the LLVM IR language which accesses the workgroup id? Frank

...o with the name you should use in your IR. Your best bet is to look at the include/llvm/IR/IntrinsicsNVVM.td file and see the definitions for each intrinsic. Then, the name mapping is just: int_foo_bar -> llvm.foo.bar() int_ prefix becomes llvm., and all underscores turn into periods. Ex: int_nvvm_read_ptx_sreg_tid_x -> llvm.nvvm.read.ptx.sreg.tid.x() On Fri, Mar 1, 2013 at 3:51 PM, Pete Couperus <pjcoup at gmail.com> wrote: > Hi Timothy, > > I'm not sure what you mean by this working for other intrinsics, but > in this case, I think you want the intrinsic name > llvm.nvvm.read.p...

...a function. class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> : NVPTXInst<(outs regclassOut:$dst), (ins), OpStr, [(set regclassOut:$dst, (IntOp))]>; def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs, int_nvvm_read_ptx_sreg_tid_x>; This method of accessing intrinsics works just fine for other intrinsics (for instance sqrt). Should I be declaring these as extern global variables? Thanks, Timothy On Fri, Mar 1, 2013 at 12:44 PM, Dmitry Mikushin <dmitry at kernelgen.org>wrote: > Timothy, > > Those calls...

...;string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> : > NVPTXInst<(outs regclassOut:$dst), (ins), > OpStr, > [(set regclassOut:$dst, (IntOp))]>; > > def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs, > int_nvvm_read_ptx_sreg_tid_x>; > > This method of accessing intrinsics works just fine for other intrinsics > (for instance sqrt). Should I be declaring these as extern global variables? > > Thanks, > > Timothy > > > On Fri, Mar 1, 2013 at 12:44 PM, Dmitry Mikushin <dmitry at kernelgen.org&g...

...lassOut, Intrinsic IntOp> : >            NVPTXInst<(outs regclassOut:$dst), (ins), >                              OpStr, >                  [(set regclassOut:$dst, (IntOp))]>; > > def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs, >    int_nvvm_read_ptx_sreg_tid_x>; > > This method of accessing intrinsics works just fine for other intrinsics > (for instance sqrt). Should I be declaring these as extern global > variables? > > Thanks, > > Timothy > > > On Fri, Mar 1, 2013 at 12:44 PM, Dmitry Mikushin > <dmitry at...

Timothy, Those calls to compute grid intrinsics are definitely wrong. In ptx code they should end up into reading special registers, rather than function calls. Try to take some working example and figure out the LLVM IR differences between it and the result of your compiler. - D. ----- Original message ----- > I've written a compiler that outputs PTX code, the result seems fairly >

I've written a compiler that outputs PTX code, the result seems fairly reasonable, but I'm not sure the intrinsics are getting compiled correctly. In addition, when I try load the module using CUDA, I get an error: CUDA_ERROR_NO_BINARY_FOR_GPU. I'm running this on a 2012 MBP with a 640M GPU. PTX Code (for a mandelbrot calculation): // // Generated by LLVM NVPTX Back-End //