search for: ptx32

...-2788-3799 #1667 -------------- next part -------------- Index: test/CodeGen/PTX/options.ll =================================================================== --- test/CodeGen/PTX/options.ll (revision 130978) +++ test/CodeGen/PTX/options.ll (working copy) @@ -1,6 +1,7 @@ ; RUN: llc < %s -march=ptx32 -mattr=ptx20 | grep ".version 2.0" ; RUN: llc < %s -march=ptx32 -mattr=ptx21 | grep ".version 2.1" ; RUN: llc < %s -march=ptx32 -mattr=ptx22 | grep ".version 2.2" +; RUN: llc < %s -march=ptx32 -mattr=ptx23 | grep ".version 2.3" ; RUN: llc < %s...

...110] [ID=1] 0x58105c0: i32 = undef [ORD=110] [ID=2] 0x58105c0: i32 = undef [ORD=110] [ID=2] The command I am using is: clang kernels/fatal_error_test.cl -O0 -include ocldef.h -include builtin_functions_ptx.cl -D__x86_64__ -ccc-host-triple ptx32 -Xclang -target-feature -Xclang +ptx23 -Xclang -target-feature -Xclang +compute20 Any ideas ? Best regards Alberto

...nsider pointer aliasing in the following kernel: __kernel void foo(__global float* a, __local float* b) { b[0] = a[0]; } If we compile this with Clang targeting PTX, the resulting LLVM IR will be: target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64" target triple = "ptx32--" define ptx_kernel void @foo(float* nocapture %a, float addrspace(4)* nocapture %b) nounwind noinline { entry: %0 = load float* %a, align 4, !tbaa !1 store float %0, float addrspace(4)* %b, align 4, !tbaa !1 ret void } !opencl.kernels = !{!0} !0 = metadata !{void (float*, float addr...

...sed an interest in using libclc as their OpenCL C library instead of developing their own. I will hope to find time over the next few weeks to add libclc support to the Clang driver. The intention is that compiling OpenCL C programs to PTX would be as easy as (something like this): clang -target ptx32 -S file.cl such that the driver would automatically locate the libclc headers, add them to the include path and pre-include the main header file. (The libclc support will of course be optional, and a -cl-stdlib= flag will be provided to allow for switching between OpenCL standard library implement...

...ibclc/ $ ./configure.py --with-llvm-config=/path/to/llvm-config && make # Then build your PTX samples $ cd ../; make build; cd build $ cmake ../ $ make --- O.K., I have a problem here. [ 12%] Compiling blur2d_kernel.cl -> blur2d_kernel.ll clang: error: no such file or directory: 'ptx32' make[2]: *** [bin/blur2d_kernel.ll] Error 1 make[1]: *** [opencl/blur2d/CMakeFiles/blur2d_kernel.ptx.dir/all] Error 2 Do you know how to fix it? Thanks. [1] https://github.com/jholewinski/llvm-ptx-samples [2] http://sites.google.com/site/justinholewinski/projects/llvm-ptx-back-end [3] http...

Hi all, Does PTX backend support llvm sitofp instruction? I failed to compile my llvm source when use llc -march=ptx32. The reason is that the source has a sitofp instruction. After i changed the instruction into uitofp manually, it passed. Thanks in advance, best, Yabin -------------- next part -------------- An HTML attachment was scrubbed... URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/201109...

...red to identify pointers [1]. I can understand that the address space attribute (.global) is not defined since there is no agreement on how to represent them in the IR. The command I use to compile is: clang kernels/parameters.cl -include ocldef.h -include builtin_functions_ptx.cl -ccc-host-triple ptx32 -S -o - where ocldef.h contains the definitions of OpenCL vector data types, builtin_functions_ptx.cl is the file provided by Justin in a previous message [2]. Is this a bug? Or the .ptr attribute is just not supported yet ? Thank you Alberto ------------------------- [1] PTX: Parallel Thread...

...s to different memory spaces are in essence different > types and cannot alias. For the kernel shown above, the resulting LLVM IR > could be: > > ; ModuleID = 'test1.cl' > target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64" > target triple = "ptx32--" > > define ptx_kernel void @foo(float* nocapture %a, float addrspace(4)* > nocapture %b) nounwind noinline { > entry: > %0 = load float* %a, align 4, !tbaa !1 > store float %0, float addrspace(4)* %b, align 4, !tbaa *!2* > ret void > } > > !opencl.kerne...

...ndef [ORD=110] [ID=2] > 0x58105c0: i32 = undef [ORD=110] [ID=2] > > The command I am using is: > > clang kernels/fatal_error_test.cl -O0 -include ocldef.h -include > builtin_functions_ptx.cl > -D__x86_64__ > -ccc-host-triple ptx32 -Xclang > -target-feature > -Xclang +ptx23 -Xclang > -target-feature > -Xclang +compute20 > > Any ideas ? > Unfortunately, this sample will not work at this time. First, the ba...

...pdate the documentation to reflect this. Also, the libclc configuration is not necessary. Only some header files are pulled in by the samples. > O.K., I have a problem here. > > [ 12%] Compiling blur2d_kernel.cl -> blur2d_kernel.ll > clang: error: no such file or directory: 'ptx32' > make[2]: *** [bin/blur2d_kernel.ll] Error 1 > make[1]: *** [opencl/blur2d/CMakeFiles/blur2d_kernel.ptx.dir/all] Error 2 > > Do you know how to fix it? Thanks. > That's a new one. I just tried a fresh clone and build, and did not experience this issue. Can you run &quot...

...32 = undef [ORD=110] [ID=2] The command I am using is: clang kernels/fatal_error_test.cl<http://fatal_error_test.cl> -O0 -include ocldef.h -include builtin_functions_ptx.cl<http://builtin_functions_ptx.cl> -D__x86_64__ -ccc-host-triple ptx32 -Xclang -target-feature -Xclang +ptx23 -Xclang -target-feature -Xclang +compute20 Any ideas ? Unfortunately, this sample will not work at this time. First, the backend does not support i8 types ye...

...are in essence different > > types and cannot alias. For the kernel shown above, the resulting LLVM IR > > could be: > > > > ; ModuleID = 'test1.cl' > > target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64" > > target triple = "ptx32--" > > > > define ptx_kernel void @foo(float* nocapture %a, float addrspace(4)* > > nocapture %b) nounwind noinline { > > entry: > > %0 = load float* %a, align 4, !tbaa !1 > > store float %0, float addrspace(4)* %b, align 4, !tbaa *!2* > > ret vo...

...ence > different > > types and cannot alias. For the kernel shown above, the resulting > LLVM IR > > could be: > > > > ; ModuleID = 'test1.cl' > > target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64" > > target triple = "ptx32--" > > > > define ptx_kernel void @foo(float* nocapture %a, float addrspace(4)* > > nocapture %b) nounwind noinline { > > entry: > > %0 = load float* %a, align 4, !tbaa !1 > > store float %0, float addrspace(4)* %b, align 4, !tbaa *!2* > > ret vo...

...;> instead of developing their own. >> >> I will hope to find time over the next few weeks to add libclc support >> to the Clang driver. The intention is that compiling OpenCL C programs >> to PTX would be as easy as (something like this): >> >> clang -target ptx32 -S file.cl >> >> such that the driver would automatically locate the libclc headers, >> add them to the include path and pre-include the main header file. >> (The libclc support will of course be optional, and a -cl-stdlib= >> flag will be provided to allow for switchi...

...ence > different > > types and cannot alias. For the kernel shown above, the resulting LLVM > IR > > could be: > > > > ; ModuleID = 'test1.cl' > > target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64" > > target triple = "ptx32--" > > > > define ptx_kernel void @foo(float* nocapture %a, float addrspace(4)* > > nocapture %b) nounwind noinline { > > entry: > > %0 = load float* %a, align 4, !tbaa !1 > > store float %0, float addrspace(4)* %b, align 4, !tbaa *!2* > > ret vo...

...their OpenCL C library > instead of developing their own. > > I will hope to find time over the next few weeks to add libclc support > to the Clang driver. The intention is that compiling OpenCL C programs > to PTX would be as easy as (something like this): > > clang -target ptx32 -S file.cl > > such that the driver would automatically locate the libclc headers, > add them to the include path and pre-include the main header file. > (The libclc support will of course be optional, and a -cl-stdlib= > flag will be provided to allow for switching between OpenCL s...

...gt; > > > > This is the testcase that I am using to verify I the max builtin > > function I am impementing > > is actually recognised. I took inspiration from the llvm-intrinsic.ll > test case. > > The command I am using to compile is: > > > > llc -march=ptx32 -mattr=+ptx22 fmax.ll > > > > The option -mattr does not seem to have any effect. > > I tried also with the ptx_device qualifier with the same outcome. > > I am using llvm from the svn repository. > > > > Bye, > > > > Alberto > > > >>&g...

...g their own. >>> >>> I will hope to find time over the next few weeks to add libclc support >>> to the Clang driver. The intention is that compiling OpenCL C programs >>> to PTX would be as easy as (something like this): >>> >>> clang -target ptx32 -S file.cl >>> >>> such that the driver would automatically locate the libclc headers, >>> add them to the include path and pre-include the main header file. >>> (The libclc support will of course be optional, and a -cl-stdlib= >>> flag will be provided...

...gt; >> > >> I will hope to find time over the next few weeks to add libclc > support > >> to the Clang driver. The intention is that compiling OpenCL C > programs > >> to PTX would be as easy as (something like this): > >> > >> clang -target ptx32 -S file.cl > >> > >> such that the driver would automatically locate the libclc headers, > >> add them to the include path and pre-include the main header file. > >> (The libclc support will of course be optional, and a -cl-stdlib= > >> flag will be prov...

...t;> I will hope to find time over the next few weeks to add libclc >> support >>>> to the Clang driver. The intention is that compiling OpenCL C >> programs >>>> to PTX would be as easy as (something like this): >>>> >>>> clang -target ptx32 -S file.cl >>>> >>>> such that the driver would automatically locate the libclc headers, >>>> add them to the include path and pre-include the main header file. >>>> (The libclc support will of course be optional, and a -cl-stdlib= >>>> fl...