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.

Maybe there are already other types like this, but it saddens me that an offline compiled code object could potentially not work properly if the application is using any of these types. Or should the runtime try to detect a problem using argument metadata?

I think it is perfectly reasonable to treat these as essentially relaxed-only atomic RMW operations and require the application to use fences or barriers if necessary. The ordering and scope are only needed if we ever need this operation to act as a non-relaxed atomic RMW.

Actually the conversions don't need non-default-rounded operations, nor are non-default-rounded arithmetic operations required by OpenCL. However, we've had requests to implement functions such as add_rtz(x,y) which computes x+y with round-to-zero rounding. Our competitors offer such functions, and we implemented them for HSAIL. So we are really trying to get back to parity with HSAIL.

One thing we've wanted for compute for quite a while now is a way to request non-default-rounded add, sub, mul, div, fma, and sqrt. Assuming we ever figure out how to represent these in the IR, ideally without falling back on intrinsics, could this approach be used to implement and minimize the mode changes for those as well?

Ping. There's quite a bit of interest in getting this exposed by clang.

This approach seems fine to me.

In D49146#1157800, @FarhanaAleen wrote:

By the way, since types are being mixed, shouldn't the summary say something like optimize fma((float)S0.x, (float)S1.x, fma((float)S0.y, (float)S1.y, S2)) --> fdot2(S0, S1, S2)? We only want this transformation if S0 and S1 are <2 x f16>.

Current pattern matching does not support float element type yet, it will be supported next.

You are right, there is a typo in the summary. It should be:
fma((f16)S0.x, (f16)S1.x fma((f16)S0.y, (f16)S1.y, (f16)z)) -> ftrunc(fdot2(S0, S1, (f32)z))

By the way, since types are being mixed, shouldn't the summary say something like optimize fma((float)S0.x, (float)S1.x, fma((float)S0.y, (float)S1.y, S2)) --> fdot2(S0, S1, S2)? We only want this transformation if S0 and S1 are <2 x f16>.

This operation only rounds a single time, and unfortunately always flushes f32 denorms. Thus this transformation should only be done when unsafe math is requested.

Thanks, looks good.

LGTM

Looks good to me.

Thanks! Looks good to me.

LGTM

LGTM, but I'd rather use '+' instead of ',' for the features.

In D39739#1045611, @ashi1 wrote:

Is first one encountered a poor design?

Looks fine to me.

Looks good to me.

In D42711#993438, @yaxunl wrote:

In D42711#992555, @arsenm wrote:

In D42711#992517, @yaxunl wrote:

I think the purpose of this patch is to get a similar usage of clang as nvptx when compiling CUDA, i.e., using cuda as OS instead of using amdhsa as OS and amdgiz as environment. This is more convenient for CUDA application developers since they just need to swap nvptx with amdgcn.

This is a frontend driver question at most. The backend shouldn't need to be aware of this

There are various places in clang where selection is done based on OS==CUDA. If we don't use that OS, we need more complex logic in clang for such choices. I can try making changes to clang to make it work, but I suspect there may be places using OS==CUDA is necessary since it may be needed before parsing the language options.

As I understand it, the option users pass to clang++ is --cuda-gpu-arch=<Arch>. Can't we arrange to generate the right triple if they use gfx900 or some other AMD target name for <Arch>?

In D42711#992503, @arsenm wrote:

You're using this just as an alias for AMDHSAOS. We shouldn't add something that behaves exactly the same

Should we expect that the last 3 arguments have any effect? Do we want to test to ensure they have the expected effects?

In D41699#990909, @b-sumner wrote:

I believe there are 20 distinct samplers. I'd be more comfortable with this if you could arrange that __translate_sampler_initializer would be called with values in [0, 19] instead of [133, 322].

I believe there are 20 distinct samplers. I'd be more comfortable with this if you could arrange that __translate_sampler_initializer would be called with values in [0, 19] instead of [133, 322].

Can this be tested?

Were you going to add min and max separately?