Address Space Qualifiers

TPC-C implements the following disjoint address spaces: __global__ and __local__. The address space qualifier may be used in variable declarations to specify the region of memory allocated to the object. The C syntax for type qualifiers is extended in TPC-C to include an address space name as a valid type qualifier.

Note

Using __global__ qualifier to allocate memory from kernel is not allowed. Use this qualifier only when reading/writing data from global memory.

A pointer to address space X can only be assigned to a pointer to the same address space X. Casting a pointer to address space X to a pointer to address space Y is illegal.

Examples:

//declares a pointer p to an int object in address space __global__

__global__ int *p;

//declares an array of 3 short128 in the __local__ address space.

__local__ short128 s[3];

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

Examples:

__local__ char f() { ... } // not allowed

   __global__ void *g() { ... } // allowed
   __local__ float64 *h() { ... } // allowed

__global__ (or _global)

The __global__ address space name is used to refer to memory objects located in the host memory from the global memory pool.

Allocating new objects in the global space is impossible. Address arithmetics with global pointers is not supported. Pointers to the global address space are allowed as arguments to functions (excluding main function).

Examples:

void main(...)

{

__global__ float64 *color; // pointer to float64 element

typedef struct {

float a[3];

int b[2];

} foo_t;

}

As tensor objects are always allocated in the global address space, do not specify __global__ qualifier for tensor keyword. The elements of a tensor object cannot be directly accessed. Built-in functions to read from and write to a tensor object are provided.

The const qualifier can also be added to the __global__ qualifier to specify a read-only memory object.

_local_ (or _local)

The _local_ address space name is used to describe variables that need to be stored in local memory. Pointers to the __local__ address space are allowed as arguments to functions (excluding main function).

Examples:

__local__ float64 buf[2] = {1.f, 2.f};

__local__ float * get_p_shift(__local__ float * p, int i)

{

return p + (i & 1);

}

void main(...)

{

// not allowed, scalar in local memory

   __local__ float a;

   // allowed, pointer to local memory, main scope

__local__ float * b = buf;

...

]