foreman
fbde6a1c77
SWDEV-184709 - support hipLaunchCooperativeKernel()
- Switch to the device library functions for GWS sync and init
Affected files ...
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#42 edit
... //depot/stg/opencl/drivers/opencl/make/hip.git/tests/src/runtimeApi/module/hipCooperativeGroup.cpp#2 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palschedcl.cpp#7 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#152 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocschedcl.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#87 edit
[ROCm/clr commit: a932af9dbc]
28 строки
589 B
C++
28 строки
589 B
C++
//
|
|
// Copyright (c) 2015 Advanced Micro Devices, Inc. All rights reserved.
|
|
//
|
|
namespace pal {
|
|
|
|
#define BLIT_KERNEL(...) #__VA_ARGS__
|
|
|
|
const char* SchedulerSourceCode = BLIT_KERNEL(
|
|
%s
|
|
\n
|
|
extern void __amd_scheduler(__global void*, __global void*, uint);
|
|
\n
|
|
__kernel void scheduler(__global void* queue, __global void* params, uint paramIdx) {
|
|
__amd_scheduler(queue, params, paramIdx);
|
|
}
|
|
\n);
|
|
|
|
const char* GwsInitSourceCode = BLIT_KERNEL(
|
|
\n
|
|
extern void __ockl_gws_init(uint nwm1, uint rid);
|
|
\n
|
|
__kernel void gwsInit(uint value) {
|
|
__ockl_gws_init(value, 0);
|
|
}
|
|
\n);
|
|
|
|
} // namespace pal
|