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 zero
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?

Looks good to me...

Do we need tests for 64-bit atomic loads? Do we need tests for cases where we get buffer instead of flat instructions?

Do we need tests of 64 bit atomic loads? What about cases where we get buffer instructions rather than flat instructions?

Do we need tests of 64-bit atomic loads? Do we need tests where we should have buffer loads instead of a flat loads?

This will break all existing implementations and is not backwards compatible. Is this extra diagnosis really worthwhile? Are developers complaining?

In D41699#971106, @yaxunl wrote:

What's the benefit of this change? Since this change will require all device libraries implementing __translate_sampler_initializer to change accordingly. We need a compelling reason.

The usual rule is to take the first weak definition encountered.

In D40514#938106, @jvesely wrote:

In D40514#937574, @b-sumner wrote:

I suppose if we have some other mechanism to ensure output modifiers are never used, then we could consider running with IEEE=0, but I think we need to continue running with IEEE=1 for compute. There may be users counting on current behavior.

The current behavior is broken for OpenCL so no one should depend on it.

IEEE mode disables output modifiers, which is good since output modifiers are not IEEE compatible and do not support output subnormal values.

Looks good to me.

In D39912#922242, @cfang wrote:

In D39912#922228, @b-sumner wrote:

Pardon my ignorance, but why isn't include/llvm/IR/IntrinsicsAMDGCN.td being updated?

We did not add new intrinsics. We just add support for new data types. In IntrinsicsAMDGCN.td, we have already defined
the data types as any_float which includes half types.

In other wordm llvm.amdgcn.image.load.v4f16 (for example) has already been declared in IntrinsicsAMDGCN.td. And this patch just needs to actually define (implement) it.

Pardon my ignorance, but why isn't include/llvm/IR/IntrinsicsAMDGCN.td being updated?

Pardon my ignorance, but why isn't include/llvm/IR/IntrinsicsAMDGPU.td being updated?

g++4.8.2 is complaining:

This passes my tests, including getting the correct answer for 0.

Looks good to me.

We could potentially update visitCEIL and visitFLOOR as well, and use the same opcode test in each, although I don't think such combinations are very likely.

LGTM. I think we can leave immediates to another patch.