search for: __local

...gt; >> Hi David, >> >> On Mon, Dec 06, 2010 at 11:14:42AM -0500, David Neto wrote: >> > What do I think your patch should look like?  It's true that the >> > diag::err_as_qualified_auto_decl is inappropriate for OpenCL when >> it's >> > the __local addres space. >> > >> > But we need to implement the semantics somehow.  Conceptually I think >> > of it as a CL source-to-source transformation that lowers >> > function-scope-local-address-space variables into a more primitive >> > form. >> >...

...t;> > >> On Mon, Dec 06, 2010 at 11:14:42AM -0500, David Neto wrote: > >> > What do I think your patch should look like?  It's true that the > >> > diag::err_as_qualified_auto_decl is inappropriate for OpenCL when > >> it's > >> > the __local addres space. > >> > > >> > But we need to implement the semantics somehow.  Conceptually I > think > >> > of it as a CL source-to-source transformation that lowers > >> > function-scope-local-address-space variables into a more primitive > >...

...ly main issue is the default address space. In OpenCL it is private, in LLVM, it is closer to global than private. > > What do I think your patch should look like? It's true that the > > diag::err_as_qualified_auto_decl is inappropriate for OpenCL when > it's > > the __local addres space. > > > > But we need to implement the semantics somehow. Conceptually I think > > of it as a CL source-to-source transformation that lowers > > function-scope-local-address-space variables into a more primitive > > form. > > > > I think I disa...

...6 -> GS, 257 -> FS). And I think 256 'standard' address spaces should be enough, but I'm happy to be proven wrong :) > What do I think your patch should look like? It's true that the > diag::err_as_qualified_auto_decl is inappropriate for OpenCL when it's > the __local addres space. > > But we need to implement the semantics somehow. Conceptually I think > of it as a CL source-to-source transformation that lowers > function-scope-local-address-space variables into a more primitive > form. > > I think I disagree that the Clang is an inappro...

...6 (and targets which lack GPU-like address > > spaces). As far as I can tell, the only real representations of > > OpenCL address spaces on such targets that we need are a way of > > distinguishing the different address spaces for alias analysis > > and a representation for __local variables allocated on the stack. > > TBAA metadata would solve the first problem, and we already have > > mechanisms in the frontend that could be used to solve the second. > > > > Which mechanisms could be used to differentiate between thread-private and > __local data...

...lack GPU-like address > > > spaces). As far as I can tell, the only real representations of > > > OpenCL address spaces on such targets that we need are a way of > > > distinguishing the different address spaces for alias analysis > > > and a representation for __local variables allocated on the stack. > > > TBAA metadata would solve the first problem, and we already have > > > mechanisms in the frontend that could be used to solve the second. > > > > > > > Which mechanisms could be used to differentiate between thread-priva...

> From: Peter Collingbourne [mailto:peter at pcc.me.uk] > Sent: 20 December 2010 20:11 > As with __local variables, it may be that "kernelness" cannot be > represented in a standard form in LLVM. For example on a CPU a > kernel function may have an additional parameter which is a pointer to > __local memory space, which would not be necessary on GPUs. Then in > fact you would u...

On Tue, Dec 21, 2010 at 07:17:40PM -0000, Anton Lokhmotov wrote: > > From: Peter Collingbourne [mailto:peter at pcc.me.uk] > > Sent: 20 December 2010 20:11 > > As with __local variables, it may be that "kernelness" cannot be > > represented in a standard form in LLVM. For example on a CPU a > > kernel function may have an additional parameter which is a pointer to > > __local memory space, which would not be necessary on GPUs. Then in > &...

...gn 4, !tbaa *!2* > > ret void > > } > > > > !opencl.kernels = !{!0} > > > > !0 = metadata !{void (float*, float addrspace(4)*)* @foo} > > *!1 = metadata !{metadata !"float$__global", metadata !3}* > > *!2 = metadata !{metadata !"float$__local", metadata !3}* > > !3 = metadata !{metadata !"omnipotent char", metadata !4} > > !4 = metadata !{metadata !"Simple C/C++ TBAA", null} > > > > Differences are bolded. Here, the TBAA pass would be able to identify > that > > the loads and st...

...nicely in the case where a CPU is > running the kernels? Does that lead to special casing or duplication > in the code generator? For example, you still have to know what > "real" calling convention to use when a kernel is running on a CPU. > (Forgive my ignorance.) As with __local variables, it may be that "kernelness" cannot be represented in a standard form in LLVM. For example on a CPU a kernel function may have an additional parameter which is a pointer to __local memory space, which would not be necessary on GPUs. Then in fact you would use a standard callin...

Indeed; I agree it is not an implementation detail as potentially valid SPIR could be almost untranslatable to valid code running on a target without dedicated workgroup-local memory. Micah: the problem can be distilled down to: __local variables in SPIR are represented as Constants (GlobalVariable : public Constant), but they are not in fact constant, for a device with no workgroup-local memory. So it is valid SPIR, as the specification stands, to manipulate __local variables as Constants in a way that is extremely difficult to...

...float %0, float addrspace(4)* %b, align 4, !tbaa *!2* > ret void > } > > !opencl.kernels = !{!0} > > !0 = metadata !{void (float*, float addrspace(4)*)* @foo} > *!1 = metadata !{metadata !"float$__global", metadata !3}* > *!2 = metadata !{metadata !"float$__local", metadata !3}* > !3 = metadata !{metadata !"omnipotent char", metadata !4} > !4 = metadata !{metadata !"Simple C/C++ TBAA", null} > > Differences are bolded. Here, the TBAA pass would be able to identify that > the loads and stores do not alias. Of course...

..._int32_base_atomics: enable #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics: enable __kernel void KMeansMapReduceAtomic(const int num_attributes, const int num_objects, __global int* delta_d ) { __local int clusterCount[256]; __local int sTemp[1]; // amd opencl needed this to be an array const unsigned int point_id = get_local_id(0); int index = 0; int i, addr; int xx = get_local_id(0); clusterCount[xx] = 0; if(get_local_id(0) == 0){ sTemp[0] = 0;...

...A, and potentially future language extensions, requires the compiler to differentiate between different types of memory. For example, OpenCL has a "__global" memory space which corresponds to globally-accessible data, and is usually off-chip memory in most GPU configurations; and a "__local" memory space which corresponds to work-group data (not accessible by work items outside of the current work group), and is usually on-chip scratchpad memory in most GPU configurations. This information is currently represented in Clang/LLVM using the addrspace() attribute on pointer types, w...

...gn 4, !tbaa *!2* > > ret void > > } > > > > !opencl.kernels = !{!0} > > > > !0 = metadata !{void (float*, float addrspace(4)*)* @foo} > > *!1 = metadata !{metadata !"float$__global", metadata !3}* > > *!2 = metadata !{metadata !"float$__local", metadata !3}* > > !3 = metadata !{metadata !"omnipotent char", metadata !4} > > !4 = metadata !{metadata !"Simple C/C++ TBAA", null} > > > > Differences are bolded. Here, the TBAA pass would be able to identify that > > the loads and stores...

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[%arrayidx1], sext...

.... Instead of calling foo(), the runtime just calls foo_kernel() which handles all of the kernel setup issues and then calls the function body itself. This removes the need to have any metadata nodes in the IR and allows the kernel function to handle any setup issues for the specific device such as __local's, id/group calculations, memory offsets, etc... without having to impact the performance of a kernel calling another kernel. Micah > -----Original Message----- > From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu] > On Behalf Of Peter Collingbourne > Sen...

Hi guys, > So it is valid SPIR, as the specification stands, to manipulate __local > variables as Constants in a way that is extremely difficult to undo. That > is, in order to transform SPIR to code that can run on a CPU, the > GlobalVariable (which is a subclass of Constant) must be replaced with a > dynamically calculated Value (which is not a subclass of constant)...

...t deciding the aliasing based on the address 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...

On 09/26/2012 08:21 PM, Villmow, Micah wrote: > It is my view that this is an implementation detail and not an issue > with the SPIR spec. As SPIR is just a representation of a program in a > portable manner, it is up to the consumer of SPIR to correctly set up > the kernels based on the devices calling convention/ABI when the SPIR > binary is loaded for that specific device. The