AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-1 - Add EMU_ENV option

Browse Source 
Fix FEAT-39125
Add EMU_ENV option. If it is ON, the building is for
emulation environment, thus some logics can be adjusted
to match emulation environment. If it is OFF, the
building is for regular environment.
Currently only ocltst will use the option. But it can
be used for other modules.

Change-Id: I54e1bc1309e82794b41fca2ae1f01f004138dced


[ROCm/clr commit: 76c6dcd558]
This commit is contained in:
taosang2
2021-11-24 21:55:27 -05:00
committed by Tao Sang
parent 8c921fc3cc
commit 6b3e3bf6c6
15 changed files with 153 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/opencl/tests/ocltst/module/runtime/OCLLDS32K.cpp
+38 -2
View File
@@ -30,12 +30,17 @@
#include <CL/cl.h>
typedef unsigned int uint32_t;
#if EMU_ENV
#define LDS_SIZE 1024
#define A_SIZE 1024
#else
#define LDS_SIZE 32768
#define A_SIZE (8 * 1024 * 1024)
#endif // EMU_ENV
#define LOCAL_WORK_SIZE 64
// We'll do a 64MB transaction
#define A_SIZE (8 * 1024 * 1024)
#define B_SIZE A_SIZE
#define C_SIZE A_SIZE
#define D_SIZE A_SIZE
@@ -46,6 +51,36 @@ typedef unsigned int uint32_t;
// 32K has 8192 elements
// 64 threads each handle 8192/64=128 values
#if EMU_ENV
static const char program_source[] = KERNEL(
__kernel void the_kernel(__global const uint *a, __global const uint *b,
__global const uint *c, __global uint *d,
__global uint *e) {
// Reduce size for the emulator
__local uint lds[256];
uint gid = get_global_id(0);
__global const uint* ta = a + 4 * gid;
__global const uint* tb = b + 4 * gid;
__global const uint* tc = c + 4 * gid;
__global uint* td = d + 4 * gid;
uint i;
for (i = 0; i < 4; ++i) lds[ta[i]] = tc[i];
barrier(CLK_LOCAL_MEM_FENCE);
for (i = 0; i < 4; ++i) td[i] = lds[tb[i]];
} __kernel void the_kernel2(__global uint* d) {
__local uint lds[8192];
uint i;
uint gid = get_global_id(0);
for (i = 0; i < 128; ++i) lds[i] = d[gid];
barrier(CLK_LOCAL_MEM_FENCE);
for (i = 0; i < 128; ++i) d[gid] = lds[i];
});
#else
static const char program_source[] = KERNEL(
__kernel void the_kernel(__global const uint *a, __global const uint *b,
__global const uint *c, __global uint *d,
@@ -73,6 +108,7 @@ static const char program_source[] = KERNEL(
for (i = 0; i < 128; ++i) d[gid] = lds[i];
});
#endif // EMU_ENV
static void fill(uint32_t *a, uint32_t *b, uint32_t *c, uint32_t *d,
uint32_t *e) {