llvm dev - Jun 2013 - [LLVMdev] How to generate constant memory for ptx code by NVPTX?

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.

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>

Thanks for reply.
I filed this bug in http://llvm.org/bugs/show_bug.cgi?id=16278.

Antony Yu



--
View this message in context:
http://llvm.1065342.n5.nabble.com/How-to-generate-constant-memory-for-ptx-code-by-NVPTX-tp58384p58416.html
Sent from the LLVM - Dev mailing list archive at Nabble.com.