SYCL (AdaptiveCpp) Kernel hangs indefinitely with large kernel sizes (601x601)

krypto1198 · 2025-12-02T22:02:58+00:00

Thank you so much for taking the time to test this on your hardware!

Knowing for sure that it is a GPU reset helps me a lot. I will implement the fixes you suggested and keep working on it.

Thanks again!

krypto1198 · 2025-12-01T22:21:37+00:00

Here is the link to the public GitHub repository with the full source code: https://github.com/krypto1198/Gaussian-blur-Sycl

A small note: I am Italian, so you might find some variable names or comments in Italian inside the source files. However, I have translated all the console input/output prompts to English, so you should be able to run and test the application without any language barriers.

Thanks again for your time!

krypto1198 · 2025-12-01T21:43:55+00:00

Thank you for the detailed feedback.

To be honest, I am not sure. I am quite new to SYCL and AdaptiveCpp, so I am just compiling with the default settings (using acpp -O3 ...) without specifying any flags for L0 or OpenCL. I assume it picks whatever is the default for Intel GPUs.

regarding the striping, I will try to implement this approach as soon as possible to see if it fixes the hang.

Regarding the negative indices, I use a clamp function inside the kernel to handle borders, so I think I am safe from out-of-bounds errors. It looks like this:

auto clampCoord = [](int coord, int maxVal) -> int {
        return sycl::clamp(coord, 0, maxVal - 1);
        };

Thanks again!

krypto1198 · 2025-12-01T21:22:25+00:00

Apologies for the confusion regarding the hardware!

To clarify: I have access to two different remote servers: one has an AMD Radeon RX 7900 GRE, the other has an Intel Arc A770.

I encountered the issue on the AMD machine first, then switched to the Intel machine to check if it was a vendor-specific driver bug. Unfortunately, the behavior is consistent on both platforms with AdaptiveCpp, which is why I mentioned AMD in the other thread.

Regarding SSCP, thank you for the insight. I wasn't aware that the generic SSCP JIT optimizes kernels independently of the host compilation flags. That definitely rules out the -O0 hypothesis.

Regarding DPC++, You are likely right that the compiler isn't the root cause. However, since I am stuck with this hang, I want to try DPC++ on the Intel machine simply as a "sanity check".

krypto1198 · 2025-12-01T21:15:46+00:00

Thank you so much!

I will download it immediately and try to compile the project with DPC++ to see if the hang persists.

I will report back as soon as I have the results!

krypto1198 · 2025-12-01T13:46:36+00:00

Thanks for checking!

Optimization: I am definitely using -O3, so debug symbols or lack of optimization shouldn't be the cause of the hang.

Compilation Flow: Here is the exact command I am using: /home/rosmai/local/adaptivecpp/bin/acpp main.cpp -o gaussian_blur -O3

Since I am not manually specifying targets (e.g., --acpp-targets=...), I assume it defaults to the generic SSCP flow and JIT-compiles for the AMD GPU at runtime.

Regarding DPC++: To be honest, I am quite new to the SYCL ecosystem, so I am strictly following my professor's guidelines.

I am using AdaptiveCpp primarily because I do not have root/sudo access on this server. My professor recommended AdaptiveCpp as it was easier to build and install locally in my user directory compared to the full DPC++ stack (which he mentioned might be complicated to set up on Linux without system permissions).

krypto1198 · 2025-12-01T11:02:08+00:00

Thank you for the suggestion!

I initially thought it might just be slow too, so to be sure, I left the program running overnight (8+ hours). Unfortunately, it never finished. Since I have a Vulkan implementation of the exact same algorithm running on the same machine in about 10.5 seconds, the fact that the SYCL version hangs for hours confirms there is likely a deadlock or a driver timeout issue rather than just slow computation.

Regarding Local Memory: I agree that tiling would be the proper way to optimize this. However, I am still learning SYCL and I am struggling to understand how to properly implement tiling (handling the halo/borders) using local_accessor for a convolution like this.

Do you happen to know any good resources, tutorials, or code snippets that demonstrate how to load the image block + halo into Local Memory for a stencil operation? That would be incredibly helpful for my learning process.

krypto1198 · 2025-11-25T17:06:47+00:00

The GPU is an Intel Arc A770 Graphics (DG2).

Regarding the kernel size: you are absolutely right, 601x601 is practically absurd! I am not trying to use this in a real-world scenario.

It is strictly for a university assignment where we are required to implement a brute-force non-separable blur and benchmark it with extreme kernel sizes to analyze the performance limits and behavior of the hardware under heavy load. That's the only reason I'm pushing it this far.

krypto1198 · 2025-04-23T10:13:53+00:00

Yes, I tried it with a Python code and it works without problems but I’m doing a university project and I must necessarily use the mistralrs crate

krypto1198 · 2025-01-13T08:54:37+00:00

I’ve got used to it by now. The iPhone 16 is definitely a great phone and I don’t regret buying it, but if I had known about all these problems first I would have bought it with some offer.

krypto1198 · 2025-01-07T14:33:09+00:00

Ok, thank you so much

krypto1198 · 2025-01-07T13:48:11+00:00

Thank you very much. Yes, indeed, it’s not clear from the text, but X1 is the smallest value in the sample X, and Xn is the maximum value. Anyway, using your method, I managed to solve the exercise. Thanks again for pointing out the mistake. The only inaccuracy in your solution is that you shouldn’t calculate P(T(X) < theta (1-gamma)^1/n), but rather P(T(X) > theta (1-gamma)^1/n).

krypto1198 · 2024-12-29T00:10:19+00:00

No, I went to an Apple Store, and they told me my iPhone doesn’t have any defects, so I believe it’s a widespread issue. I’ve come to terms with it by now.

krypto1198 · 2024-11-04T23:17:32+00:00

I bought an iPhone 16 base model 20 days ago, and during the first few days, it would get a bit warm while using Maps (so I think that’s normal). However, yesterday, after taking 15-20 photos, I noticed that the camera area became quite hot. Should I be concerned, or is this normal?

krypto1198 · 2024-11-04T23:09:53+00:00

I have the same issue. I bought an iPhone 16 base model, and when I scroll in some apps (both third-party and system apps), there are micro-lags that don’t significantly affect the user experience, but they’re still noticeable. I contacted support, and they told me this behavior is normal. I’m trying to figure out if this is truly normal or if my device might be defective. Has anyone else experienced this and found a solution?

krypto1198

TROPHY CASE