SYCL atomics are slow on MI250

[mgrabban]

SYCL atomics run slow on AMD MI250.
Multiple workloads show performance degradation compared to MI100, and after looking into details, we find that the reason for slower performance is kernels that have atomics.

Can you please test SYCL atomics on AMD MI250?

FYI I tested in -

• Ubuntu 20.04.4 LTS
• rocm-5.2.3

$ clang++ --version
clang version 16.0.0 (https://github.com/intel/llvm 10d978425267ea8543587048606c4d8550dbff95)
Target: x86_64-unknown-linux-gnu
...

Note that I have been seeing slow performance on MI250 ever since I started using it, which was at least 4 months back. So it is probably not due to open source sycl version or rocm version.

Please let me know if you need more information. Thanks.

[zjin-lcf]

Are the HIP assembly codes for the atomics the same ?

[al42and]

Are you using -munsafe-fp-atomics flag and/or unsafeAtomicAdd/atomicAddNoRet intrinsics?

[mgrabban]

Are the HIP assembly codes for the atomics the same ?

Can you please tell me how to generate the HIP assembly code from either SYCL source or clang++ compiled exe?

[mgrabban]

Are you using -munsafe-fp-atomics flag and/or unsafeAtomicAdd/atomicAddNoRet intrinsics?

No, I am not using any of these flags in either MI100 or MI250. I did not know about them. Can you link to their docs?

[al42and]

Can you link to their docs?

https://github.com/ROCm-Developer-Tools/HIP/blob/develop/docs/markdown/hip_kernel_language.md#unsafe-floating-point-atomic-rmw-operations

And ORNL website has a good description of what exactly is "safe"/"unsafe" in the context of MI250X: https://docs.olcf.ornl.gov/systems/crusher_quick_start_guide.html, search for "unsafe-fp-atomics".

Setting -munsafe-fp-atomics should, theoretically, use the appropriate intrinsic automatically.

[mgrabban]

Ok I want to add more details here

• We have HIP versions of our workloads and do not see them running slower on MI250 compared to MI100.
• Only for SYCL versions, kernels with atomics run slower on MI250. So for some reason, clang++ is generating slower machine code. Note that we have atomics involving uint32_t, float as well as double and all runs slower so data type is probably not the reason for performance issue.

I tried using -munsafe-fp-atomics on MI250 and -mno-unsafe-fp-atomics on MI100: I do not see any impact on performance but I do see these warnings

MI250
clang-16: warning: argument unused during compilation: '-munsafe-fp-atomics' [-Wunused-command-line-argument]

MI100
clang-16: warning: argument unused during compilation: '-mno-unsafe-fp-atomics' [-Wunused-command-line-argument]

For my test, I am using this code

#include <sycl/sycl.hpp>

#include <iostream>
#include <chrono>

#define atomicAdd(x, y) (sycl::atomic_ref<float, sycl::memory_order::relaxed, sycl::memory_scope::work_group, sycl::access::address_space::global_space>(*(x)) +=(y))

const int factor     = 4;
const int num_points = 10240 * 1 * 256;
const int num_wgs    = 10240 * factor;
const int wg_size    = 256 / factor;

const int tm_size    = 32;
const int num_iters  = 1;

// add 1.0f to array element if it is in my 'team'
void kernel(float* data, sycl::nd_item<1> item)
{
    size_t tid = item.get_global_linear_id();
    size_t sid = tid / tm_size;

    for (size_t i = 0; i < tm_size; i++) {
        atomicAdd(data + (sid * tm_size + i), 1.0f);
    }
}

int main()
{
    try {
    sycl::device dev = sycl::device(sycl::default_selector_v);
    sycl::queue q(dev);

    float* data_d = sycl::malloc_device<float>(num_points, q);
    float* data_h = sycl::malloc_host<float>(  num_points, q);

    // warm-up run:
    for (int iter = 0; iter < 1; iter++) {
        q.memset(data_d, 0.0, num_points * sizeof(float));
        q.wait();

        q.parallel_for(
            sycl::nd_range<1>(num_points, wg_size),
            [=](sycl::nd_item<1> item) {
                kernel(data_d, item);
            }
        );
        q.wait_and_throw();
    }

    // timed run:
    std::chrono::time_point<std::chrono::steady_clock> time11;
    std::chrono::time_point<std::chrono::steady_clock> time12;
    double duration = 0.0;

    for (int iter = 0; iter < num_iters; iter++) {
        q.memset(data_d, 0.0, num_points * sizeof(float));
        q.wait();

        time11 = std::chrono::steady_clock::now();
        q.parallel_for(
            sycl::nd_range<1>(num_points, wg_size),
            [=](sycl::nd_item<1> item) {
                kernel(data_d, item);
            }
        );
        q.wait_and_throw();
        time12 = std::chrono::steady_clock::now();
        duration += std::chrono::duration<double, std::micro>(time12 - time11).count();
    }

    q.memcpy(data_h, data_d, num_points * sizeof(float));
    q.wait();

    for (size_t i = 0; i < num_points; i++) {
        if (data_h[i] != (float)tm_size) {
            std::cerr << "mismatch found at index " << i << std::endl;
            return 3;
        }
    }
    std::cout << "Success!\n";
    std::cout << "num_points: " << num_points << std::endl;
    std::cout << "num_wgs: " << num_wgs << std::endl;
    std::cout << "wg_size: " << wg_size << std::endl;
    std::cout << "tm_size: " << tm_size << std::endl;
    std::cout << "Total time for " << num_iters << " iterations: " << duration << " us\n";

    } catch (sycl::exception const &exc) {
		std::cerr << exc.what() << "Exception caught at file:" << __FILE__
				  << ", line:" << __LINE__ << std::endl;
		return 1;
	} catch (std::exception& e) {
		std::cerr << "Uncaught exception: " << e.what() << std::endl;
        return 2;
	}

    return 0;
}

This code runs 136 times faster on MI100 compared to MI250. The work-group size is optimized for MI250 and I see MI100 performance can be increased further by tuning the work-group size.

[al42and]

Interesting. When I compile your code for both MI100 and MI250 ($ clang++ -fsycl -O3 -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx908 test.cpp -save-temps), I get almost identical code in both cases. But the gfx90a version has additional buffer_wbl2/buffer_invl2 instructions around the global_atomic_cmpswap.

[al42and]

Ah, my bad. It seems in the HIP backend, the native floating point atomics are not implemented at all, so it will always be a CAS loop. Then indeed, no effect is to be had from -munsafe-fp-atomics and the like.

Upd: BTW, I'd expect that your MI100 HIP code can benefit from -munsafe-fp-atomics if it's not somehow already enabled. Will only affect float atomics, but still something.

[mgrabban]

Interesting. When I compile your code for both MI100 and MI250 ($ clang++ -fsycl -O3 -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend --offload-arch=gfx908 test.cpp -save-temps), I get almost identical code in both cases. But the gfx90a version has additional buffer_wbl2/buffer_invl2 instructions around the global_atomic_cmpswap.

So this could be a reason for slow performance?

[mgrabban]

Ah, my bad. It seems in the HIP backend, the native floating point atomics are not implemented at all, so it will always be a CAS loop. Then indeed, no effect is to be had from -munsafe-fp-atomics and the like.

Upd: BTW, I'd expect that your MI100 HIP code can benefit from -munsafe-fp-atomics if it's not somehow already enabled. Will only affect float atomics, but still something.

Yes we will use -munsafe-fp-atomics for HIP versions of our workloads. Thanks for the suggestion.

[al42and]

So this could be a reason for slow performance?

No idea. But that's the only difference I see in the generated assembly between gfx908 and gfx90a. Can't find any documentation either (I'm not affiliated with AMD), but it looks like a cache invalidation.

A somewhat bigger problem, performance-wise, is that CAS-loops are used for float atomics of MI250 at all. I don't know whether oneAPI/libclc team has any timeline for lifting this restriction.

[mgrabban]

Ok I will keep pursuing this.
Thanks again for all the help :-)

[zjin-lcf]

#4471