From d33c1a06cfd6cfd7648693a24ca1511451a718f5 Mon Sep 17 00:00:00 2001 From: Ruili Ji Date: Tue, 18 Apr 2023 14:34:48 +0800 Subject: [PATCH] SWDEV-391103 - fix ocltst segment fault and update device info for gfx1103 Change-Id: I7a45ef98fd03113b148178fed3b64b919ee4eb6d --- .../ocltst/module/perf/OCLPerfPipeCopySpeed.cpp | 16 ++++++++-------- .../ocltst/module/runtime/OCLDeviceQueries.cpp | 1 + 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/opencl/tests/ocltst/module/perf/OCLPerfPipeCopySpeed.cpp b/opencl/tests/ocltst/module/perf/OCLPerfPipeCopySpeed.cpp index 7224ff8400..940a3c46b5 100644 --- a/opencl/tests/ocltst/module/perf/OCLPerfPipeCopySpeed.cpp +++ b/opencl/tests/ocltst/module/perf/OCLPerfPipeCopySpeed.cpp @@ -149,9 +149,9 @@ const static char * strKernel = {\n int gid = get_global_id(0);\n local reserve_id_t resId;\n - resId = sub_group_reserve_write_pipe(outPipe, get_local_size(0));\n + resId = sub_group_reserve_write_pipe(outPipe, get_sub_group_size());\n if (is_valid_reserve_id(resId)) {\n - write_pipe(outPipe, resId, get_local_id(0), &inBuf[gid]);\n + write_pipe(outPipe, resId, get_sub_group_local_id(), &inBuf[gid]);\n sub_group_commit_write_pipe(outPipe, resId);\n }\n }\n @@ -159,14 +159,14 @@ const static char * strKernel = kernel __attribute__((reqd_work_group_size(64,1,1))) void copyPipe_sg(read_only pipe DATA_TYPE inPipe, write_only pipe DATA_TYPE outPipe)\n {\n local reserve_id_t resId;\n - resId = sub_group_reserve_read_pipe(inPipe, get_local_size(0));\n + resId = sub_group_reserve_read_pipe(inPipe, get_sub_group_size());\n if (is_valid_reserve_id(resId)) {\n DATA_TYPE tmp;\n - read_pipe(inPipe, resId, get_local_id(0), &tmp);\n + read_pipe(inPipe, resId, get_sub_group_local_id(), &tmp);\n sub_group_commit_read_pipe(inPipe, resId);\n - resId = sub_group_reserve_write_pipe(outPipe, get_local_size(0));\n + resId = sub_group_reserve_write_pipe(outPipe, get_sub_group_size());\n if (is_valid_reserve_id(resId)) {\n - write_pipe(outPipe, resId, get_local_id(0), &tmp);\n + write_pipe(outPipe, resId, get_sub_group_local_id(), &tmp);\n sub_group_commit_write_pipe(outPipe, resId);\n }\n }\n @@ -176,10 +176,10 @@ const static char * strKernel = {\n int gid = get_global_id(0);\n local reserve_id_t resId;\n - resId = sub_group_reserve_read_pipe(inPipe, get_local_size(0));\n + resId = sub_group_reserve_read_pipe(inPipe, get_sub_group_size());\n if (is_valid_reserve_id(resId)) {\n DATA_TYPE tmp;\n - read_pipe(inPipe, resId, get_local_id(0), &tmp);\n + read_pipe(inPipe, resId, get_sub_group_local_id(), &outBuf[gid]);\n sub_group_commit_read_pipe(inPipe, resId);\n outBuf[gid] = tmp;\n }\n diff --git a/opencl/tests/ocltst/module/runtime/OCLDeviceQueries.cpp b/opencl/tests/ocltst/module/runtime/OCLDeviceQueries.cpp index e78a300db5..d5b91bde2f 100644 --- a/opencl/tests/ocltst/module/runtime/OCLDeviceQueries.cpp +++ b/opencl/tests/ocltst/module/runtime/OCLDeviceQueries.cpp @@ -116,6 +116,7 @@ static const AMDDeviceInfo DeviceInfo[] = { /* Navi31*/ { "gfx1100", "gfx1100", 4, 32, 1, 256, 64 * Ki, 32, 11, 0 }, /* Navi32*/ { "gfx1101", "gfx1101", 4, 32, 1, 256, 64 * Ki, 32, 11, 0 }, /* Navi33*/ { "gfx1102", "gfx1102", 4, 32, 1, 256, 64 * Ki, 32, 11, 0 }, + /* Phoenix */ { "gfx1103", "gfx1103", 4, 32, 1, 256, 64 * Ki, 32, 11, 0 }, }; const int DeviceInfoSize = sizeof(DeviceInfo) / sizeof(AMDDeviceInfo);