Skip to content

Commit

Permalink
Add attribute settings to random access
Browse files Browse the repository at this point in the history
  • Loading branch information
Mellich committed May 13, 2020
1 parent 77f0bc5 commit 8881ee1
Show file tree
Hide file tree
Showing 2 changed files with 26 additions and 1 deletion.
Original file line number Diff line number Diff line change
@@ -0,0 +1,19 @@

global_memory_name = "HBM"

def generate_attributes(num_replications, num_global_memory_banks=32):
"""
Generates the kernel attributes for the global memory. They specify in which
global memory the buffer is located. The buffers will be placed using a
round robin scheme using the available global memory banks and the number of
replications that should be generated (e.g. if a global memory contains multiple banks)
@param num_replications Number okernel replications
@param num_global_memory_banks Number of global memory banks that should be used for generation
@return Array of strings that contain the attributes for every kernel
"""
global_memory_names = [ "%s%d" % (global_memory_name, i) for i in range(num_global_memory_banks)]
return [ "__attribute__((buffer_location(\"%s\")))"
% (global_memory_names[i % num_global_memory_banks])
for i in range(num_replications)]
8 changes: 7 additions & 1 deletion RandomAccess/src/device/random_access_kernels_single.cl
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,12 @@ Constant used to update the pseudo random number
*/
#define POLY 7

/* PY_CODE_GEN
try:
kernel_param_attributes = generate_attributes(num_replications)
except:
kernel_param_attributes = ["" for i in range(num_replications)]
*/

// PY_CODE_GEN block_start [replace(local_variables=locals()) for i in range(num_replications)]

Expand All @@ -42,7 +48,7 @@ to the kernel.
*/
__attribute__((max_global_work_dim(0)))
__kernel
void accessMemory_/*PY_CODE_GEN i*/(__global DEVICE_DATA_TYPE_UNSIGNED volatile * restrict data,
void accessMemory_/*PY_CODE_GEN i*/(__global /*PY_CODE_GEN kernel_param_attributes[i]*/ DEVICE_DATA_TYPE_UNSIGNED volatile * restrict data,
const DEVICE_DATA_TYPE_UNSIGNED m,
const DEVICE_DATA_TYPE_UNSIGNED data_chunk,
const uint kernel_number) {
Expand Down

0 comments on commit 8881ee1

Please sign in to comment.