diff --git a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h index 33623f5881..e1dc54f270 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h +++ b/projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h @@ -91,7 +91,6 @@ extern "C" __device__ __attribute__((const)) double __ocml_cvtrtn_f64_u64(__hip_ extern "C" __device__ __attribute__((const)) double __ocml_cvtrtp_f64_u64(__hip_uint64_t); extern "C" __device__ __attribute__((const)) double __ocml_cvtrtz_f64_u64(__hip_uint64_t); -extern "C" __device__ __attribute__((convergent)) void __ockl_gws_init(uint nwm1, uint rid); extern "C" __device__ __attribute__((convergent)) void __ockl_gws_barrier(uint nwm1, uint rid); extern "C" __device__ __attribute__((const)) __hip_uint32_t __ockl_lane_u32(); diff --git a/projects/clr/rocclr/device/blitcl.cpp b/projects/clr/rocclr/device/blitcl.cpp index 19a23a1cf6..b2f6aebe7b 100644 --- a/projects/clr/rocclr/device/blitcl.cpp +++ b/projects/clr/rocclr/device/blitcl.cpp @@ -47,8 +47,6 @@ const char* BlitLinearSourceCode = BLIT_KERNELS( extern void __ockl_dm_init_v1(ulong, ulong, uint, uint); - extern void __ockl_gws_init(uint nwm1, uint rid); - __kernel void __amd_rocclr_fillBufferAligned(__global void* buf, __constant uchar* pattern, uint pattern_size, uint alignment, ulong end_ptr, uint next_chunk, uint workgroup_size) { @@ -182,7 +180,7 @@ const char* HipExtraSourceCode = BLIT_KERNELS( __ockl_dm_init_v1(heap_to_initialize, initial_blocks, heap_size, number_of_initial_blocks); } - __kernel void __amd_rocclr_gwsInit(uint value) { __ockl_gws_init(value, 0); }); + __kernel void __amd_rocclr_gwsInit(uint value) { __builtin_amdgcn_ds_gws_init(value, 0); }); const char* HipExtraSourceCodeNoGWS = BLIT_KERNELS( __kernel void __amd_rocclr_streamOpsWrite(__global uint* ptrInt, __global ulong* ptrUlong,