From 447beeb00b31b560a003863a54f932b6499ec75c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Mon, 15 Dec 2025 15:56:14 +0000 Subject: [PATCH] Replace usages of __ockl_gws_init with __builtin_amdgcn_ds_gws_init (#2235) --- .../clr/hipamd/include/hip/amd_detail/device_library_decls.h | 1 - projects/clr/rocclr/device/blitcl.cpp | 4 +--- 2 files changed, 1 insertion(+), 4 deletions(-) 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,