__LOCAL(3clc)

📦
version 1.0~svn33624 · 3 versions
, in ubuntu jammy, ubuntu noble

✨ __LOCAL(3clc) ✨

okay bestie, buckle up, we're diving into the OpenCL Manual and this one's all about __local — the address space qualifier that's basically the group-chat of memory. 💅

🌸 NAME 🌸

__local — Address Space Qualifier.

also known as: __local, local

🦋 DESCRIPTION 🦋

so OpenCL is out here implementing these disjoint address spaces, and there are exactly four of them, no more no less:

  • __global
  • __local
  • __constant
  • __private

The address space qualifier can be used in variable declarations to specify which region of memory is used to allocate the object. The C syntax for type qualifiers gets extended in OpenCL so that an address space name counts as a valid type qualifier. So if an object's type is qualified by an address space name, the object gets allocated in that specified address name — and if it's not qualified? it lands in the generic address space. periodt.

here's the cute part: the address space names without the __ prefix — that's global, local, constant and private — may be substituted for the corresponding address space names with the __ prefix. so you can be lazy and it's fine, slay.

The generic address space name for arguments to a function in a program, or for local variables of a function, is __private. And this is a hard rule, no cap: all function arguments shall be in the __private address space.

functionQualifiers(3clc) function arguments declared to be a pointer of a type can only point to one of these address spaces:

  • __global
  • __local
  • __constant

A pointer to address space A can only be assigned to a pointer to the same address space A. Casting a pointer to address space A into a pointer to address space B? illegal. not "discouraged", not "a vibe", illegal. 💀

Function arguments of type otherDataTypes(3clc), otherDataTypes(3clc), otherDataTypes(3clc), otherDataTypes(3clc), otherDataTypes(3clc), and otherDataTypes(3clc) refer to image memory objects allocated in the __global address space.

🎀 NOTES 🎀

The __local or local address space name describes variables that need to be allocated in local memory and are shared by all work-items of a work-group. it's giving shared google doc energy fr fr.

Pointers to the __local address space are allowed as arguments to functions — including __kernel functions. 📌

Variables allocated in the __local address space can also be defined inside a __kernel function scope. iconic.

✨ Example ✨

There is no generic address space name for program scope variables. All program scope variables must be declared in the __constant address space. For example:

// declares a pointer p in the
__private address space that // points
to an int object in address space
__global __global int *p;

// declares an array of 4 floats in the
__private address space.  float x[4];

There is no address space for function return values. Using an address space qualifier in a function return type declaration will generate a compilation error — unless the return type is declared as a pointer type and the qualifier is used on the points-to address space. iykyk.

For example:

__private int f() { ... } // should
generate an error __local int *f() {
... } // allowed __local int *
__private f() { ... }; // should
generate an error.

// Examples of variables allocated in
the __local address space // inside a
__kernel function

__kernel void my_func(...)  { __local
float      a;       // A single float
allocated // in local address space

__local float      b[10];   // An array
of 10 floats // allocated in local
address space.  if (...)  { // example
of variable in __local address space
but not // declared at __kernel
function scope.  __local float c;
// not allowed.  } }

Variables allocated in the __local address space inside a __kernel function cannot be initialized. the way that is such a trap bestie 😭

__kernel void my_func(...)
{
local float a = 1;    // not allowed __local float b; b = 1;                // allowed
}
NOTE: Variables allocated in the __local address space inside a __kernel function are allocated for each work-group executing the kernel.

⭐ SPECIFICATION ⭐

OpenCL Specification[1]

👑 SEE ALSO 👑

💜 AUTHORS 💜

The Khronos Group

Copyright © 2007-2011y The Khronos Group Inc.

Permission is hereby granted, free of charge, to any person obtaining a copy of this software and/or associated documentation files (the "Materials"), to deal in the Materials without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Materials, and to permit persons to whom the Materials are furnished to do so, subject to the following condition:

the above copyright notice and this permission notice shall be included in all copies or substantial portions of the Materials.

🧚 NOTES 🧚

  1. OpenCL Specification, page 222, section 6.5 - Address Space Qualifiers

okay that's the whole page bestie, you ate that 🍓 the OpenCL Manual says hi, this came from The Khronos Group and it's dated 01/14/2021, filed under __LOCAL(3clc) ✨💖🌸

📦 where this page comes from