[LLVMdev] How to generate constant memory for ptx code by NVPTX?

Justin Holewinski justin.holewinski at gmail.com
Fri Jun 7 16:42:54 PDT 2013


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( ...){    ...  }
>
> 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 one_f = 0f3F800000;
> .entry test( ...) { ... }
>
> one_f is placed in global memory and load it by ld.global.f32
> But I want one_f to be placed in constant memory and load it by
> ld.const.f32. Is it any way to achieve this ??
>
> Thanks in advance.
> Antony Yu
>
>
>
> --
> View this message in context:
> http://llvm.1065342.n5.nabble.com/How-to-generate-constant-memory-for-ptx-code-by-NVPTX-tp58384.html
> Sent from the LLVM - Dev mailing list archive at Nabble.com.
> _______________________________________________
> LLVM Developers mailing list
> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev
>



-- 

Thanks,

Justin Holewinski
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20130607/ab123901/attachment.html>


More information about the llvm-dev mailing list