Page MenuHomePhabricator

mariusz-sikora-at-amd (Mariusz Sikora)
User

Projects

User does not belong to any projects.

User Details

User Since
Mar 2 2022, 12:56 AM (65 w, 1 d)

Recent Activity

Apr 20 2023

mariusz-sikora-at-amd added inline comments to D148796: [AMDGPU][GFX908] Add builtin support for global add atomic f16/f32.
Apr 20 2023, 6:44 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd added reviewers for D148796: [AMDGPU][GFX908] Add builtin support for global add atomic f16/f32: arsenm, rampitec, gandhi21299, foad.

This is a proposal how we could extend global fadd atomic builtins support for gfx908

Apr 20 2023, 5:42 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D148796: [AMDGPU][GFX908] Add builtin support for global add atomic f16/f32.
Apr 20 2023, 5:39 AM · Restricted Project, Restricted Project, Restricted Project

Mar 30 2023

mariusz-sikora-at-amd updated the diff for D147159: [AMDGPU][GFX940] Simplify definitions of GLOBAL_ATOMIC.

Add tests for gfx908 with unsupported instructions.

Mar 30 2023, 2:04 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D147159: [AMDGPU][GFX940] Simplify definitions of GLOBAL_ATOMIC.
Mar 30 2023, 2:03 AM · Restricted Project, Restricted Project

Mar 29 2023

mariusz-sikora-at-amd added reviewers for D147159: [AMDGPU][GFX940] Simplify definitions of GLOBAL_ATOMIC: rampitec, foad, piotr.
Mar 29 2023, 7:27 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D147159: [AMDGPU][GFX940] Simplify definitions of GLOBAL_ATOMIC.
Mar 29 2023, 7:25 AM · Restricted Project, Restricted Project

Mar 24 2023

mariusz-sikora-at-amd committed rG69061f96275c: [AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16 (authored by mariusz-sikora-at-amd).
[AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16
Mar 24 2023, 8:28 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd closed D146808: [AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16.
Mar 24 2023, 8:28 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd added reviewers for D146808: [AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16: foad, rampitec, piotr, gandhi21299.
Mar 24 2023, 5:57 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D146808: [AMDGPU] Add clang builtin for __builtin_amdgcn_ds_atomic_fadd_v2f16.
Mar 24 2023, 5:54 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd committed rGea064ee2a3bd: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions (authored by mariusz-sikora-at-amd).
[AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions
Mar 24 2023, 5:12 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd closed D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.
Mar 24 2023, 5:12 AM · Restricted Project, Restricted Project, Restricted Project

Mar 23 2023

mariusz-sikora-at-amd updated the diff for D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.

Put back Real Instructions under SubtargetPredicate associated with gfx generation.

Mar 23 2023, 1:50 PM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.
Mar 23 2023, 12:59 PM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.
Mar 23 2023, 7:55 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.
Mar 23 2023, 7:45 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd updated the diff for D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.

Changes after review.

Mar 23 2023, 6:35 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.
Mar 23 2023, 4:47 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd added reviewers for D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions: rampitec, foad, piotr.
Mar 23 2023, 2:01 AM · Restricted Project, Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions.
Mar 23 2023, 1:58 AM · Restricted Project, Restricted Project, Restricted Project

Jan 23 2023

mariusz-sikora-at-amd committed rG848978db4ccb: [AMDGPU][NFC] Apply new naming convention for feature fmacf64 (authored by mariusz-sikora-at-amd).
[AMDGPU][NFC] Apply new naming convention for feature fmacf64
Jan 23 2023, 10:00 PM · Restricted Project, Restricted Project
mariusz-sikora-at-amd closed D142329: [AMDGPU][NFC] Apply new naming convention for feature fmacf64.
Jan 23 2023, 10:00 PM · Restricted Project, Restricted Project

Jan 22 2023

mariusz-sikora-at-amd added reviewers for D142329: [AMDGPU][NFC] Apply new naming convention for feature fmacf64: arsenm, rampitec.
Jan 22 2023, 10:18 PM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D142329: [AMDGPU][NFC] Apply new naming convention for feature fmacf64.
Jan 22 2023, 10:16 PM · Restricted Project, Restricted Project

Jan 19 2023

mariusz-sikora-at-amd added inline comments to D142017: [AMDGPU] Add feature predicate for v_fmac_f64 instruction.
Jan 19 2023, 6:59 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd committed rGbf4140c2768e: [AMDGPU] Add feature predicate for v_fmac_f64 instruction (authored by mariusz-sikora-at-amd).
[AMDGPU] Add feature predicate for v_fmac_f64 instruction
Jan 19 2023, 4:07 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd closed D142017: [AMDGPU] Add feature predicate for v_fmac_f64 instruction.
Jan 19 2023, 4:07 AM · Restricted Project, Restricted Project

Jan 18 2023

mariusz-sikora-at-amd added reviewers for D142017: [AMDGPU] Add feature predicate for v_fmac_f64 instruction: foad, rampitec.
Jan 18 2023, 6:57 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D142017: [AMDGPU] Add feature predicate for v_fmac_f64 instruction.
Jan 18 2023, 6:55 AM · Restricted Project, Restricted Project

Jan 17 2023

mariusz-sikora-at-amd committed rGe87934415499: [AMDGPU] v_fmac_f64 encoding tests for gfx940 (authored by mariusz-sikora-at-amd).
[AMDGPU] v_fmac_f64 encoding tests for gfx940
Jan 17 2023, 1:08 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd closed D141857: [AMDGPU] v_fmac_f64 encoding tests for gfx940.
Jan 17 2023, 1:08 AM · Restricted Project, Restricted Project

Jan 16 2023

mariusz-sikora-at-amd added reviewers for D141857: [AMDGPU] v_fmac_f64 encoding tests for gfx940: foad, rampitec.
Jan 16 2023, 8:41 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D141857: [AMDGPU] v_fmac_f64 encoding tests for gfx940.
Jan 16 2023, 8:39 AM · Restricted Project, Restricted Project

Nov 28 2022

mariusz-sikora-at-amd added reviewers for D138786: [AMDGPU] Add llvm.amdgcn.raw.atomic.buffer.load intrinsic to support OpAtomicLoad lowering: dstuttard, foad, piotr, nhaehnle.
Nov 28 2022, 3:51 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D138786: [AMDGPU] Add llvm.amdgcn.raw.atomic.buffer.load intrinsic to support OpAtomicLoad lowering.
Nov 28 2022, 3:48 AM · Restricted Project, Restricted Project

Oct 18 2022

mariusz-sikora-at-amd added a comment to D135110: [NFC] [HLSL] Move common metadata to LLVMFrontend.

I don't quite know why, but it seems like this new library breaks tests of llvm-config; if I start out with an empty build directory and run ninja check-llvm (or ninja llvm-test-depends), then this library doesn't get built, and llvm-config prints llvm-config: error: missing: $HOME/code/llvm-project/llvm/build/lib/libLLVMFrontendHLSL.a - and that breaks tools/llvm-config/booleans.test and tools/llvm-config/system-libs.test.

Oct 18 2022, 2:33 AM · Restricted Project, Restricted Project, Restricted Project

Oct 12 2022

mariusz-sikora-at-amd abandoned D135437: [Tests] Require zlib for zstd tests.

Closing this review. This issue will be fixed in review : https://reviews.llvm.org/D135744

Oct 12 2022, 3:46 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd accepted D135744: [llvm-objcopy] Support --decompress-debug-sections when zlib is disabled.

This change is fixing my issue. Thanks !
I will close my review.

Oct 12 2022, 3:40 AM · Restricted Project, Restricted Project

Oct 10 2022

mariusz-sikora-at-amd added a comment to D135437: [Tests] Require zlib for zstd tests.

Sorry for not providing more details.
These tests may fail if zlib is installed on the system, but front end is explicitly disabling LLVM_ENABLE_ZLIB in cmake, then LIT tests will fail with information that LLVM_ENABLE_ZLIB has been disabled.

Oct 10 2022, 6:36 AM · Restricted Project, Restricted Project

Oct 7 2022

mariusz-sikora-at-amd requested review of D135437: [Tests] Require zlib for zstd tests.
Oct 7 2022, 3:28 AM · Restricted Project, Restricted Project

May 4 2022

mariusz-sikora-at-amd committed rG2417de2758c3: [AMDGPU] Use d16 flag for image.sample instructions (authored by mariusz-sikora-at-amd).
[AMDGPU] Use d16 flag for image.sample instructions
May 4 2022, 9:30 PM · Restricted Project, Restricted Project
mariusz-sikora-at-amd closed D124232: [AMDGPU] Use d16 flag for image.sample instructions.
May 4 2022, 9:30 PM · Restricted Project, Restricted Project

Apr 28 2022

mariusz-sikora-at-amd added a comment to D124232: [AMDGPU] Use d16 flag for image.sample instructions.

I wanted to add this combine before, but I don’t think there is a way to add d16 to an instruction without potentially breaking the code.
The reason is, when an image_sample has the d16 flag enabled, it will use f32→f16 truncation or i32→i16 truncation, depending on the texture format in the descriptor.

Combining image_sample+fptrunc to image_sample d16 works fine for float textures, but I assume we don’t know at compile time if a texture is an integer or float texture.
The application may interpret stored values as float and does an fptrunc, but the texture is actually defined as an integer texture, so the hardware uses an integer trunc instead, giving different results.

Are you aware of any examples that would fail with this patch? At least in Vulkan, the image format must match the sampled type.

I see, I didn’t know the type on the intrinsic always matches the type of the texture. Sorry for the noise.
It would be nice to have the same combine with integer textures.

Apr 28 2022, 3:52 AM · Restricted Project, Restricted Project

Apr 27 2022

mariusz-sikora-at-amd updated the diff for D124232: [AMDGPU] Use d16 flag for image.sample instructions.

Compare pointers of OldIntr and InstToRemove instead of using isIdenticalTo.

Apr 27 2022, 4:11 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D124232: [AMDGPU] Use d16 flag for image.sample instructions.
Apr 27 2022, 3:42 AM · Restricted Project, Restricted Project

Apr 26 2022

mariusz-sikora-at-amd updated the diff for D124232: [AMDGPU] Use d16 flag for image.sample instructions.

In last version (which was merged and reverted) I made a mistake at the end
of modifyIntrinsicCall function where I was doing operations on removed
instruction (InstToReplace).

Apr 26 2022, 9:41 PM · Restricted Project, Restricted Project
mariusz-sikora-at-amd reopened D124232: [AMDGPU] Use d16 flag for image.sample instructions.

This change was merged to main and later reverted because of memory-sanitizer failing tests. I'm re-opening this review to publish updated version of the change.

Apr 26 2022, 9:37 PM · Restricted Project, Restricted Project

Apr 25 2022

mariusz-sikora-at-amd added a comment to D124232: [AMDGPU] Use d16 flag for image.sample instructions.

I don't have commit access, can anyone merge this patch for me, please?

Apr 25 2022, 4:46 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D124217: [AMDGPU] Allow finer grain control of an unaligned access speed.
Apr 25 2022, 1:01 AM · Restricted Project, Restricted Project

Apr 22 2022

mariusz-sikora-at-amd updated the diff for D124232: [AMDGPU] Use d16 flag for image.sample instructions.

Changes after review

Apr 22 2022, 10:11 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd added inline comments to D124232: [AMDGPU] Use d16 flag for image.sample instructions.
Apr 22 2022, 10:10 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd added reviewers for D124232: [AMDGPU] Use d16 flag for image.sample instructions: piotr, arsenm, foad, dstuttard.
Apr 22 2022, 12:41 AM · Restricted Project, Restricted Project
mariusz-sikora-at-amd requested review of D124232: [AMDGPU] Use d16 flag for image.sample instructions.
Apr 22 2022, 12:15 AM · Restricted Project, Restricted Project