In D148906#4321445, @foad wrote:

I think we need to similarly remove a wavefrontsize64 function if compiling in wave32 mode.

I think this is a reasonable request but it's an extension to what this pass currently does.

• Currently the pass removes functions that rely on features not supported by the -mcpu= GPU.
• The extension would be to remove functions that rely on features not supported by the -mcpu= GPU as modified by any -mattr= options.

Does that make sense? If so can we please implement that in a generic way instead of having a special hack for wave32/wave4 features?

I think we need to similarly remove a wavefrontsize64 function if compiling in wave32 mode.

How confident are we in our ability to strip the convergent attribute off of functions that don't need it? Seems like this could cause performance regressions.

I think a function with wavefrontsize32 is incompatible with a wavefrontsize64 compilation and similarly for a wavefrontsize64 function with a wavefrontsize32 compilation and we should take the usual steps for incompatible functions.

In D148906#4287110, @foad wrote:

In D148906#4286903, @arsenm wrote:

In D148906#4286691, @foad wrote:

Currently, wavfrontsize32 is being appended by device-libs to some functions for gfx9 targets.

Why? That seems wrong.

It is wrong but that’s how the library is currently structured. All code coexists in the same IR, and there are some wave32 only functions

But if you're compiling for gfx900 then this pass should remove any functions that are marked with +wavefrontsize32. It should not just remove the +wavefrontsize32 from the function.

In D148796#4284504, @rampitec wrote:

We used to support it that way and decided just not doing it. It is very hard to explain why a supported atomic results in error. Someone who really needs it can use intrinsic.

In D147408#4243403, @foad wrote:

Scalar branches may be the most expensive aspect of this algorithm

If not-taken conditional branches are cheap then we could do something like this. It only has one taken branch, when we have finished handling all the active lanes.

  // Inclusive plus-scan v0 into v1. Also leaves the result of the plus-reduction in s3.
  s_mov s0, exec
  s_mov s3, 0 // accumulator
// repeat this section 32 or 64 times:
  s_ff1 s1, s0 // find lowest remaining active lane
  s_cmp_eq s1, -1
  s_cbranch_scc1 end
  s_bitset0 s0, s1
  v_readlane s2, v0, s1
  s_add s3, s2
  v_writelane v1, s3, s1
// end of repeated section
end:

No objection here.

It might be better to loop over only active lanes, found using ctz or clz builtins.

In D146523#4210894, @foad wrote:

In D146523#4210843, @b-sumner wrote:

FWIW, there is no desire to read from inactive lanes. The loop is supposed to only be reading from, and writing to, lanes that were active before the for loop is executed by a select single lane.

Then I'm back to not understanding what this convergent copy is for. I'd need to see a more complete example.

FWIW, there is no desire to read from inactive lanes. The loop is supposed to only be reading from, and writing to, lanes that were active before the for loop is executed by a select single lane.

In D145343#4170305, @yaxunl wrote:

In D145343#4170250, @arsenm wrote:

I think exposing whether or not the flag was used is weird/broken, as is including _OPTION in the name. Should just define to whether it's enabled or not

I agree. @b-sumner What do you think?

My current understanding is the c-p will go into already forked clang-16, but not to rocm 5.4. So rocm device-libs will be accompanied by the older clang-16 w/o this and stay compatible. Someone building from scratch will use latest clang-16 and staging device-libs with this change. Do you think this will work?

In D142507#4127382, @rampitec wrote:

In D142507#4127374, @aaronmondal wrote:

I think unless conflicts arise creating an issue similar to this https://github.com/llvm/llvm-project/issues/60600 with the cherry-pick line set to this commit should be enough. (See also https://llvm.org/docs/GitHub.html).

I believe it will need D142407 to be cherry-picked as well to apply cleanly. Otherwise I do not expect conflicts. So the c-p need to go into release/16.x, right?
Let's wait for @b-sumner first anyway, he is maintaining device-lib.

In D143643#4115531, @arsenm wrote:

In D143643#4115510, @b-sumner wrote:

The device library is using the builtin, so I assume the build will fail when this change lands.

This doesn’t change the builtin at all

Also, I see no median intrinsic described in the LLVM programming guide.

The device library is using the builtin, so I assume the build will fail when this change lands.

No objections here...

__builtin_sqrtf does not produce a correctly rounded result. I don't recommend this change.

Thank you!

We're in a lose-lose situation here. No matter what we do, someone will complain. But we're are definitely on more solid footing by going for correctness.

Different functions providing different behaviors can be handled at link time like any other function, instead of the same functions providing different behaviors per translation unit and requires cloning. The current scheme transfers complexity from the device library build system into the driver and user binaries

In D130096#3850473, @arsenm wrote:

In D130096#3850472, @jhuber6 wrote:

I don't like the fact that we need to have two different kinds of control constants, one per-TU and others per-link job. I'm wondering how difficult it would be to make the fast versions of the math calls use different entry points. That way we could handle this in the math header wrappers.

That's really how the C linkage model wants you to handle this. I also would like to have FP value tracking optimizations take care of the special cases in the library code

Also, we may want to use uppercase for other purposes in the future.

I don't particularly see a need for this. I am not opposed to a "did you mean" in the error diagnostic.

In D134355#3809294, @yaxunl wrote:

In D134355#3807435, @cfang wrote:

LGTM

Should the module flag name be amdgpu_code_object_version or amdhsa_code_object_version?

Good question.

@b-sumner Does code object version affects PAL? Thanks.

In D132140#3732337, @yaxunl wrote:

revised by Brian's comments

Following existing naming, it might make sense to rename "rtn_b32" --> "rtn" and "rtn_b64" --> "rtnl".

In D131276#3702934, @arsenm wrote:

I also think we're still missing a module flag to indicate the code object version

In D99071#3652434, @nickdesaulniers wrote:

Are declarations for @llvm.amdgcn.is.shared supposed to be showing up in my IR, even when my target triple is _not_ amdgcn-amd?

The "is_" prefix is not helpful. I suggest either "uses_dynamic_stack" or just "dynamic_stack".

Thanks @bcahoon. If there were another available approach that always worked, we would use it.

This should include updates to https://llvm.org/docs/AMDGPUUsage.html#kernel-descriptor if not planned for a different patch.

t-tye should review this too.

I think this is too focused. There are other image_sample_lz intrinsics and all of them potentially need to be replaced if following this approach.

Looks fine to me!

This change goes along with 1194b9cdda30d and should probably have been landed with it.

In D120150#3332613, @foad wrote:

Ran the test over the entire range of double floating-point.

Just curious: how long does it take to test all 2^64 inputs?

In D118229#3291464, @arsenm wrote:

In D118229#3291454, @yaxunl wrote:

In D118229#3290769, @arsenm wrote:

In D118229#3290747, @yaxunl wrote:

In D118229#3290659, @arsenm wrote:

In D118229#3290658, @kzhuravl wrote:

This change introduces the regression in OpenCL conformance test: basic - kernel_memory_alignment_local. Does it require any corresponding runtime changes?

Is clang correctly emitting the align attribute on all these arguments?

clang does not do anything special for alignment of pointer type kernel arg. It assumes the pointee alignment is default 1. https://godbolt.org/z/xs195rKoW

Question is: What OpenCL spec says about kernel arg pointee alignment? @b-sumner @Anastasia

It should be the ABI type alignment as was used before

I doubt OpenCL spec requires alignment of pointer-type kernel argument. I suspect it is part of our own undocumented ABI. If we implement this in clang, it probably goes to TargetABIInfo.

Yes it does, it requires natural alignment for all types which is what that conformance test is checking.

Can we agree to drop the macro in LLVM 15 and note that in a comment or elsewhere?

In D114957#3166936, @foad wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I don't know anything about OpenCL macros. Is it good enough to put this in AMDGPUTargetInfo::getTargetDefines:

if (Opts.OpenCL)
  Builder.defineMacro("__amdgcn_bvh_use_vec3__");

Does it need tests, documentation, etc?

In D114957#3166882, @yaxunl wrote:

In D114957#3166861, @arsenm wrote:

In D114957#3166858, @yaxunl wrote:

In D114957#3166817, @foad wrote:

This is a flag-day change to the signatures of the LLVM intrinsics and the OpenCL builtins. Is that OK?

This breaks users' code. If we have to do this, at least let clang emit a pre-defined macro e.g. __amdgcn_bvh_use_vec3__=1 so that users can make their code work before and after the change.

I do not think it's worth introducing a macro for this. Are there actually C users of these builtins?

Yes we have users who use these clang builtins. We have received quite a few complaints about making breaking API changes without a way to detect them in the program.

I believe that for the purposes of detecting lane 0 mbcnt_lo is sufficient.

Two approaches for limiting the stores to lane 0 of each wave:

1. Write 1 to exec mask, store, and write -1 to exec mask. This works since the exec mask at the start of the wave when this happens is -1
2. Check for lane == 0 and branch. The lane can be computed by a) wave64: builtin_amdgcn_mbcnt_hi(~0u, builtin_amdgcn_mbcnt_lo(~0u, 0u)) b) wave32: __builtin_amdgcn_mbcnt_lo(~0u, 0u)

This looks good to me.

That is the safest thing to do. However, I'm now hearing some thinking that memory likely to be targeted by f.p. atomics is likely to be cached. But again, better to be safe at least until we're clear this will always be the case.

Thanks. LGTM.

That sequence will give a very accurate result as long as overflow and underflow is avoided. LGTM.

Should this also be IntrConvergent?

Looks good to me.

In D82818#2141981, @arsenm wrote:

In D82818#2141973, @kzhuravl wrote:

Do we also want to remove it from v2 metadata?

Probably, but I looked briefly and didn't actually see the direct equivalent

LGTM

In D79580#2026301, @arsenm wrote:

In D79580#2026263, @rampitec wrote:

In D79580#2026107, @arsenm wrote:

In D79580#2025743, @rampitec wrote:

In D79580#2025725, @arsenm wrote:

In D79580#2025685, @rampitec wrote:

How do we get these address spaces? I'd rather error out.

You can use attribute((address_space)) or hand write IR. We should not error on valid IR. Someday a frontend may wish to track special information with a custom address space number, which we can handle as a global alias similar to how x86 accepts any arbitrary address space

It sounds dangerous. An unknown address space has unknown semantics. How could we assume anything about it?

It doesn't have unknown semantics, it has target defined semantics. We can interpret them however we want. At a minimum any pointer computation on the values should work, even if load and store don't select

Is there any particular motivation to add this support?

Just generally making a best effort attempt to handle any IR that passes the verifier. We already try to handle these in a variety of other places the same way, and theoretically someone could want to use it.

In D77910#2015225, @arsenm wrote:

In D77910#1989163, @b-sumner wrote:

In D77910#1988807, @arsenm wrote:

In D77910#1981828, @b-sumner wrote:

In D77910#1981429, @arsenm wrote:

In D77910#1976171, @b-sumner wrote:

I don't think we can guarantee this is or will be supported on all devices. The language runtime makes this decision.

We don't need to worry about theoretical devices. We should know the properties of the driver from -amdhsa, -amdpal, -mesa3d

It takes more than support in the ISA for some features. The OpenCL driver may not want to support a given optional feature, e.g. images. I'm not opposed to defaults, but if the driver chooses to not support images, it needs to be able to prevent __IMAGE_SUPPORT__ from being defined. Conformance will fail if the runtime and compiler are not consistent.

The driver details should be captured by the the triple. If some weird driver decided to do something different, we would need to add a new triple for it. We don't have such a driver, so I don't see why worry about it. It's possible to work around with undef and redef in an implicitly included header. We need to fix properties of the driver based on the target to have perfectly matching offline compilation

I don't see anywhere in the triple talking about driver specific details, unless you would use the environment? That seems like overkill to me. But again, I'm not opposed to defaults, and as long as the driver can override them, this should be OK.

The OS is the driver. It doesn't need to specifically encode these details; the OS should encode properties of the driver environment. Anything using -amdhsa should be reporting image support

In D77910#1988807, @arsenm wrote:

In D77910#1981828, @b-sumner wrote:

In D77910#1981429, @arsenm wrote:

In D77910#1976171, @b-sumner wrote:

I don't think we can guarantee this is or will be supported on all devices. The language runtime makes this decision.

We don't need to worry about theoretical devices. We should know the properties of the driver from -amdhsa, -amdpal, -mesa3d

It takes more than support in the ISA for some features. The OpenCL driver may not want to support a given optional feature, e.g. images. I'm not opposed to defaults, but if the driver chooses to not support images, it needs to be able to prevent __IMAGE_SUPPORT__ from being defined. Conformance will fail if the runtime and compiler are not consistent.

The driver details should be captured by the the triple. If some weird driver decided to do something different, we would need to add a new triple for it. We don't have such a driver, so I don't see why worry about it. It's possible to work around with undef and redef in an implicitly included header. We need to fix properties of the driver based on the target to have perfectly matching offline compilation

In D77910#1981429, @arsenm wrote:

In D77910#1976171, @b-sumner wrote:

I don't think we can guarantee this is or will be supported on all devices. The language runtime makes this decision.

We don't need to worry about theoretical devices. We should know the properties of the driver from -amdhsa, -amdpal, -mesa3d

In my opinion, for on-line compile for OpenCL, the platform is responsible for setting __OPENCL_VERSION__. Also, it should be the platform's choice as to how to respond the image support query and how __IMAGE_SUPPORT__ is set. For offline compile, it doesn't seem unreasonable to ask the developer to set these.

I don't think we can guarantee this is or will be supported on all devices. The language runtime makes this decision.

In addition to predefining __ATOMIC_RELAXED, etc., clang also predefines __OPENCL_MEMORY_SCOPE_WORK_ITEM and friends. So it doesn't really seem unreasonable for clang to also predefine its known syncscopes, and to require the argument to be one of those integers.

LGTM

Please go ahead and update to a string for the scope.

Thanks. This looks fine to me.

I'm not aware of any valid explicit uses of this intrinsic by library or developer code, so this seems fine to me.

I recall we agreed that conformance tests using mixed types were broken, so this change should be OK. Hopefully this will not affect users.

LGTM

Should this be looking forward to also handling OpenCL, which does require vector support?

I am OK with proceeding here.

I am OK with proceeding here.

In D71358#1780083, @arsenm wrote:

In D71358#1780059, @arsenm wrote:

In D71358#1780038, @rampitec wrote:

How is it now achieved that we do not support denormals by default before GFX9?

This is a frontend decision. clang emits the attribute based on the subtarget preference

This is also how it always worked. The backend uniformly treated denormals as off by default

This looks OK to me, although tuning on correctly rounded division any time denorms are enabled is not actually required by OpenCL.

In D66197#1629784, @jdoerfert wrote:

Do we really need these to be "amdgpu" specific?

Looks fine to me.

Looks fine to me.

In D66197#1629436, @arsenm wrote:

In D66197#1629348, @b-sumner wrote:

Looks fine to me. Thanks!

I wonder if is.local should be is.shared because that's what getreg calls this

Looks fine to me. Thanks!

Looks fine to me. Thanks!

In D62739#1543438, @arsenm wrote:

In D62739#1536428, @b-sumner wrote:

We need to communicate with anyone generating IR to ensure this is being generated before we change the default. clang is only one of those generators. This change will also need to be documented in the usage document.

The planned change is to make the backend more conservative, so it shouldn't break other frontends

We need to communicate with anyone generating IR to ensure this is being generated before we change the default. clang is only one of those generators. This change will also need to be documented in the usage document.

Looks good to me.