search for: __global

Dear all, on porting blender 2.69 to FreeBSD -CURRENT (which by default uses libc++ and clang 3.3), I ran into an issue with the cycles render engine implementation, which apparently defines several OpenCL-specific keywords, such as __global, __local, ... The issue about defining compiler-specific keywords aside, the compilation failed due to an issue in /usr/include/c++/v1/__locale, which defines a private member static locale& __global(); in the locale class (see also http://llvm.org/viewvc/llvm-project/libcxx/trunk/include/...

...PrU3AS2cz, while in clang's opencl-c.h[2], printf does not have the overload attribute: int printf(__constant const char* st, ...); (and it is different from the standard, which is printf(restrict __constant char *, ...)) I try the following code: #include <opencl-c.h> __kernel void vadd(__global const int* a, __global const int* b, __global int* c) { printf("aaaaa"); } and get a printf that is not mangled: %call = tail call spir_func i32 (i8 addrspace(2)*, ...) @printf(i8 addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)* @.str, i64 0, i64 0)) with the f...

...SPIR 2.0 spec[1], the name of OpenCL builtins are mangled. > > > > However, when I compile OpenCl code with Clang 3.9 with the > "spir64-unknown-unknown" target, Clang generates IR without mangling the > builtins, e.g. for: > > > > __kernel void input_zip_int(__global int *in0) { > > *in0 = get_global_id(0); > > } > > > > clang generates: > > > > define spir_kernel void @input_zip_int(i32 addrspace(1)* nocapture %in0) > local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 > !kernel_arg_type !5 !ke...

...PrU3AS2cz, while in clang's opencl-c.h[2], printf does not have the overload attribute: int printf(__constant const char* st, ...); (and it is different from the standard, which is printf(restrict __constant char *, ...)) I try the following code: #include <opencl-c.h> __kernel void vadd(__global const int* a, __global const int* b, __global int* c) { printf("aaaaa"); } and get a printf that is not mangled: %call = tail call spir_func i32 (i8 addrspace(2)*, ...) @printf(i8 addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)* @.str, i64 0, i64 0)) with the f...

...1 > #20 0x00007ffff6b5d179 in clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) () from /opt/buildbot/lib/libOpenCL.so.1 > #21 0x00007ffff6b1282c in (anonymous namespace)::compile_llvm (llvm_ctx=..., > source="\n__kernel void test_fn(__local float *sSharedStorage, __global float *srcValues, __global uint *offsets, __global float *destBuffer, uint alignmentOffset )\n{\n int tid = get_global_id( 0 );\n sSha"..., headers=..., name="input.cl", triple="r600--", processor="verde", opts="", > address_spaces=..., opti...

The `Metadata`/`Value` split (PR21532) landed in r223802 -- at least, the C++ side of it. This was a rocky day, but I suppose that's what I get for failing to stage the change in smaller pieces. As of r223916 (lldb), I'm not aware of any remaining (in-tree) breakage, so if I've missed some problem in the sea of buildbot errors, please flag me down. I'll follow up soon with

...vate variable to point to the one copy of that struct. The pointer is returned by a system-defined intrinsic function dependent on the current work item. (The system knows what work groups are in flight, which is why you need a system-defined intrinsic.) So a kernel function like this: void foo(__global int*A) { __local int vint; __local int *vpint; __local int const *vcpint; __local int volatile vvint; int a = A[0]; vint = a; vvint = a; int a2 = vint; int va2 = vvint; barrier(CLK_LOCAL_MEM_FENCE); A[0] = a2 + va2; } is translated to this, which does pass through...

...007ffff6b5d179 in clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) () from /opt/buildbot/lib/libOpenCL.so.1 >>> #21 0x00007ffff6b1282c in (anonymous namespace)::compile_llvm (llvm_ctx=..., >>> source="\n__kernel void test_fn(__local float *sSharedStorage, __global float *srcValues, __global uint *offsets, __global float *destBuffer, uint alignmentOffset )\n{\n int tid = get_global_id( 0 );\n sSha"..., headers=..., name="input.cl", triple="r600--", processor="verde", opts="", >>> address_spaces=.....

Hi all, According to the SPIR 2.0 spec[1], the name of OpenCL builtins are mangled. However, when I compile OpenCl code with Clang 3.9 with the "spir64-unknown-unknown" target, Clang generates IR without mangling the builtins, e.g. for: __kernel void input_zip_int(__global int *in0) { *in0 = get_global_id(0); } clang generates: define spir_kernel void @input_zip_int(i32 addrspace(1)* nocapture %in0) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { entry: %call = ta...

...f that > struct. The pointer is returned by a system-defined intrinsic > function dependent on the current work item. (The system knows what > work groups are in flight, which is why you need a system-defined > intrinsic.) > > So a kernel function like this: > > void foo(__global int*A) { > __local int vint; > __local int *vpint; > __local int const *vcpint; > __local int volatile vvint; > int a = A[0]; > vint = a; > vvint = a; > int a2 = vint; > int va2 = vvint; > barrier(CLK_LOCAL_MEM_FENCE); > A[0] = a2 +...

...lover + an opencl kernel 1 and 2 have been tested on a kepler card, 3 has been tested with pocl. My goal for this week is to get the tgsi backend to produce code which I can copy and paste into 2 and then have it working on a kepler card. The test program looks like this: __kernel void test_kern(__global uint *vals, __global uint *buf) { uint id = get_global_id(0); buf[32 * id] -= vals[id]; } The llvm ir looks like this: bin/clang -x cl -c -emit-llvm -target tgsi-- -include /usr/share/pocl/include/_kernel.h -o ~/foo.ir -x cl -S ~/foo.cl ; ModuleID = '/home/hans/foo.cl' target data...

Hi everybody, I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks. Compiling the Histogram64.cl program I get a several backend errors. I isolated one of them in the following kernel program: __kernel void kernel_function(__global int *input) { __local char localArray[16]; for(unsigned int index = 0; index < 16; ++index) localArray[index] = 0; input[0] = localArray[get_local_id(0)]; } fatal error: error in backend: Cannot select: 0x5810cc0: i32,ch = load 0x57fa148, 0x5810ac0, 0x58105c0<LD...

...clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) () from /opt/buildbot/lib/libOpenCL.so.1 > >>>> #21 0x00007ffff6b1282c in (anonymous namespace)::compile_llvm (llvm_ctx=..., > >>>> source="\n__kernel void test_fn(__local float *sSharedStorage, __global float *srcValues, __global uint *offsets, __global float *destBuffer, uint alignmentOffset )\n{\n int tid = get_global_id( 0 );\n sSha"..., headers=..., name="input.cl", triple="r600--", processor="verde", opts="", > >>>> address_s...

...lerInstance::ExecuteAction(clang::FrontendAction&) () from /opt/buildbot/lib/libOpenCL.so.1 >>>>>>> #21 0x00007ffff6b1282c in (anonymous namespace)::compile_llvm (llvm_ctx=..., >>>>>>> source="\n__kernel void test_fn(__local float *sSharedStorage, __global float *srcValues, __global uint *offsets, __global float *destBuffer, uint alignmentOffset )\n{\n int tid = get_global_id( 0 );\n sSha"..., headers=..., name="input.cl", triple="r600--", processor="verde", opts="", >>>>>>> add...

Hi, When trying to compile a trivial opencl kernel such as: __kernel void if_eq(__global int * out, int arg0, int arg1){ out[0] = arg0==arg1?0:1; } Clang generates IR like: %1 = icmp eq i32 %arg0, %arg1 %. = zext i1 %1 to i32 This eventually crashes ISel on R600. Attached patch adds a selector so it will compile. Regards, Jon Pry jonpry at gmail.com -------------- ne...

...ss space should happen > before considering type informations, so the TBAA extension would not be a > general solution. > I could see this going either way. You could argue that address space information is a part of the type system (e.g. a "__local float*" cannot alias a "__global float *") for the source-level aliasing information, and not part of the type system for target-dependent address space checks (e.g. in PTX, address space 1 cannot alias address space 2). I do not have a strong preference either way. Though I agree from a performance standpoint that checking...

Hi everybody, I am trying to produce ptx code starting from OpenCL C. I am experiencing a problem concerning pointer parameters. Here follows an example: kernel void function(__global float* parameter1) {} NVIDIA NVCC Compiler: .entry function( .param .u32 *.ptr* .global .align 4 function_param_0 ) { ret; } CLANG + LLVM PTX backend // (skipping builtin functions definitions) .entry function (.param .b32 __param_1) // @function { // BB#0: //...

Dear Developers, I compiled a OpenCL kernel before (on Nov. last year) like __kernel void g(__global float* array) { array[get_global_id(0)] = 1; } with libclc, which would originally use the instrinsics like llvm.r600.read.local.size.x(). I executed the generated object file with one version of the hsa-runtime [1] provided by Mr. Stellard, when there was more than one workgroup, the output of...

...gmail.com>wrote: > Hi everybody, > > I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks. > Compiling the Histogram64.cl program I get a several backend errors. > > I isolated one of them in the following kernel program: > > __kernel void kernel_function(__global int *input) { > __local char localArray[16]; > for(unsigned int index = 0; index < 16; ++index) > localArray[index] = 0; > input[0] = localArray[get_local_id(0)]; > } > > fatal error: error in backend: Cannot select: > 0x5810cc0: i32,ch = load 0x57fa148...

....com<mailto:alberto.magni86 at gmail.com>> wrote: Hi everybody, I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks. Compiling the Histogram64.cl program I get a several backend errors. I isolated one of them in the following kernel program: __kernel void kernel_function(__global int *input) { __local char localArray[16]; for(unsigned int index = 0; index < 16; ++index) localArray[index] = 0; input[0] = localArray[get_local_id(0)]; } fatal error: error in backend: Cannot select: 0x5810cc0: i32,ch = load 0x57fa148, 0x5810ac0, 0x58105c0<LD1[%arr...