Databases
StackOverflow2013
Postsdbo
POWhat is the best way to implement a small lookup table in an OpenCL Kernel
 

Note that there are some explanatory texts on larger screens.

plurals
  1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
    primarykey
    data
    text
    <p>In my kernel it is necessary to make a large number of random accesses to a small lookup table (only 8 32-bit integers). Each kernel has a unique lookup table. Below is a simplified version of the kernel to illustrate how the lookup table is used.</p> <pre><code>__kernel void some_kernel( __global uint* global_table, __global uint* X, __global uint* Y) { size_t gsi = get_global_size(0); size_t gid = get_global_id(0); __private uint LUT[8]; // 8 words of of global_table is copied to LUT // Y is assigned a value from the lookup table based on the current value of X for (size_t i = 0; i &lt; n; i++) { Y[i*gsi+gid] = LUT[X[i*gsi+gid]]; } } </code></pre> <p>Because of the small size, I am getting the best performance by keeping the table in the __private memory space. However, because of the random nature in which the lookup table is accessed, there is still a large performance hit. With the lookup table code removed (replaced with a simple arithmetic operation, for example), although the kernel would provide the wrong answer, the performance improves by a factor of over 3.</p> <p>Is there a better way? Have I overlooked some OpenCL feature that provides efficient random access for very small chunks of memory? Could there be an efficient solution using vector types?</p> <p>[edit] Note, that the maximum value of X is 7, but the maximum value of Y is as large as 2^32-1. In other words, all the bits of the lookup table are being used, so it cannot be packed into a smaller representation.</p>
    singulars
    1. PO
      singulars
      1. PTAnswer
    1. This table or related slice is empty.
    1. PTQuestion
    1. USDustin
    1. USDustin
    plurals
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PO
      singulars
      1. PTAnswer
    2. PO
      singulars
      1. PTAnswer
    1. VO
      singulars
      1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
      1. This table or related slice is empty.
      1. VTUpMod
    2. VO
      singulars
      1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
      1. USprunge
      1. VTFavorite
    3. VO
      singulars
      1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
      1. This table or related slice is empty.
      1. VTUpMod
    1. COJust to make sure I am understanding correctly is LUT[] and X[] unique for every single work item?
      singulars
      1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
      1. USAdam S.
    2. COHave you tried using `__constant` memory for your look-up tables? GPUs often implement separate caches and memory access pathways for constant memory, for speeding up things like shared look-up tables.
      singulars
      1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
      1. USuser57368
    3. COI did try passing global_table as __constant, but for some reason, it didn't help the performance. Each kernel operates on an independent section of global_table, X, and Y (depending on the thread ID).
      singulars
      1. POWhat is the best way to implement a small lookup table in an OpenCL Kernel
      1. USDustin
 

Querying!

 
Guidance
A row detail

Detail views are divided into sections. All the information in the data section comes from columns in the selected row. The other sections display data from other, related rows.

Related data can be related in a to-one or a to-many fashion. Captions of data related in a to-many fashion link to a list view showing a filtered view of the table.

Try moving around until you find a non-empty to-many entry and click on the label to get to one. You can move back to the root by clicking on the database name in the header.

SQuiL has stopped working due to an internal error.

If you are curious you may find further information in the browser console, which is accessible through the devtools (F12).

Reload