Commit 2c2b93ee authored by Bartlomiej Wroblewski's avatar Bartlomiej Wroblewski
Browse files

Remove nop from inline asm

parent 5469c613
...@@ -977,7 +977,6 @@ __device__ void amd_direct_load_global_to_lds(const T* global_base_ptr, ...@@ -977,7 +977,6 @@ __device__ void amd_direct_load_global_to_lds(const T* global_base_ptr,
auto const lds_ptr_sgpr = auto const lds_ptr_sgpr =
__builtin_amdgcn_readfirstlane((reinterpret_cast<uintptr_t>(lds_ptr))); __builtin_amdgcn_readfirstlane((reinterpret_cast<uintptr_t>(lds_ptr)));
asm volatile("s_mov_b32 m0, %0; \n\t" asm volatile("s_mov_b32 m0, %0; \n\t"
"s_nop 0;\n\t"
"buffer_load_dword %1, %2, 0 offen lds;\n\t" ::"s"(lds_ptr_sgpr), "buffer_load_dword %1, %2, 0 offen lds;\n\t" ::"s"(lds_ptr_sgpr),
"v"(global_offset_bytes), "v"(global_offset_bytes),
"s"(src_resource)); "s"(src_resource));
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment