[SYCL][NFCI] Use DeviceConfigFile instead of sycl_aspects metadata

[jzc]

The sycl_aspects metadata kept track of the aspect names and their enum values as defined by sycl/include/sycl/info/aspects.def. This metadata was populated by the frontend by looking for the declaration of the SYCL aspects enum. The CleanupSYCLMetadata pass then deletes this metadata after device compilation.

llvm/clang/lib/CodeGen/BackendUtil.cpp

Lines 1156 to 1158 in f6e73e8

	// Remove SYCL metadata added by the frontend, like sycl_aspects
	// Note, this pass should be at the end of the pipeline
	MPM.addPass(CleanupSYCLMetadataPass());

However, some processing in sycl-post-link that will be added in #12727 needs the sycl_aspects metadata (which at that point, the metadata is already deleted by CleanupSYCLMetadata). To fix this, in this PR, the aspect enum values were added to DeviceConfigFile.td. By adding this, this makes other uses of the sycl_aspects data redundant, so the places where sycl_aspects were needed were replace by uses of the device config file. (essentially just SYCLPropagateAspectsUsage.cpp.)

Since the aspect values in the device config file are duplicated from aspects.def, a test was added in device_config_file_aspects.cpp to ensure that these values stay in sync with aspects.def. Additionally, some other tests had to be updated due to the sycl_aspects metadata no longer being present.

[premanandrao]

What happens for this testcase after your changes?

[jzc]

No error or note will be emitted, I removed the code that would do so otherwise

[premanandrao]

Is it not possible for the user to redefine the aspect enum anymore? I understand that you have removed the diagnostic for it, but is it safe to do so?

[jzc]

From my understanding, yes, it is safe to remove this diagnostic, as it is not relevant anymore.

The only purpose that I found of the [[__sycl_detail__::sycl_type(aspect)]] attribute was to communicate what declaration is the aspect enum declartion (stored in AspectsEnumDecl member), which then uses that to add the sycl_aspects metadata listing the aspect names and values. But since this information is now provided by the device config file and I've removed all AspectsEnumDecl, having two enum declarations with [[__sycl_detail__::sycl_type(aspect)]] attached to them doesn't do anything, so therefore, it is safe to remove the diagnostic.

[AlexeySachkov]

However, some processing in sycl-post-link that will be added in #12727 needs the sycl_aspects metadata (which at that point, the metadata is already deleted by CleanupSYCLMetadata).

To clarify this point for others: in #12727 we need to compare values listed in device config file (strings representing aspects) with aspects metadata attached to kernels (integers that come from aspect enumeration) and it is impossible to achieve without knowing the mapping between the two.

Whilst this PR does simplify our implementation by removing special handling of this enum from FE, the PR also defeats the original purpose of this special handling entirely. That purpose is to decouple the compiler from the runtime. See #5892

Even though there is a test case suggested that ensures that we config file contents matches the SYCL headers to prevent hard to debug errors, it is impossible to use the same compiler with a different SYCL headers without updating the enum values, which could be an ABI break.

One of the main ideas of the existing approach is that we have a generic infrastructure in the compiler, which would work with different implementations in headers by providing an interface in form of [[__sycl_detail__::sycl_type(aspect)]] to let the compiler know which exact enumeration is special and contains those aspect values. Adding an attribute to header files is a much smaller modification than potentially changing values of enumerators.

Tagging @tahonermann here for awareness, because it impacts our upstreaming effort.

Perhaps what we want is to encode strings right away into metadata we generate from [[sycl::device_has()]] attribute and corresponding property. This shouldn't be too complex with the property (even though it will require compile-time strings concatenation, see virtual functions design doc proposal for a prototype) and I don't think that it would change processing of properties by the passes much, but it may require more tricks around [[sycl::device_has()]] attribute. So it is once again important to hear from FE folks and folks who work on upstreaming efforts.

[jzc]

@AlexeySachkov Thanks for the context, that makes sense. If we decide not to go through with this change, then in terms of follow up for #12727, directly emitting string attributes make sense. Perhaps a simpler solution is instead of CleanupSYCLMetadata removing all the metadata keeping track of the mapping, it keep only the aspects used. For the general use case, this would reduce the mapping down to something 1-3 associations if using things like half, double, atomic64, etc., instead of the ~60 or so total associations that usually are there.

[steffenlarsen]

@AlexeySachkov Thanks for the context, that makes sense. If we decide not to go through with this change, then in terms of follow up for #12727, directly emitting string attributes make sense. Perhaps a simpler solution is instead of CleanupSYCLMetadata removing all the metadata keeping track of the mapping, it keep only the aspects used. For the general use case, this would reduce the mapping down to something 1-3 associations if using things like half, double, atomic64, etc., instead of the ~60 or so total associations that usually are there.

From what I remember (@asudarsa can correct me if I am misremembering) we needed the cleanup pass because the linker would append the aspect lists, which could cause the aspect lists to be extremely long, despite it being filled with only a handful unique values. Since we didn't previously need them at- or after linking, cleanup pre-linking seemed like a good solution. Reducing the aspects to only those used might help, but we don't avoid the lists growing during linking. Maybe there is a better solution to prevent the aspects lists growing in linking? Can we change the representation of this list in a way that the lists don't get merged?

[asudarsa]

Do we expect this to grow? Otherwise, we can drop the SmallVector and the loop. Just a nit. Thanks

[asudarsa]

Hi @jzc

Looks like a good change. I had a couple of high level questions:

1. Is this an NFCI? IIUC Metadata generation is impacted here.
2. Is it not possible to just push the cleanupSYCLMetadata pass to run at the end of sycl-post-link?

Thanks

[jzc]

@asudarsa

Is this an NFCI? IIUC Metadata generation is impacted here.

That is true, although my intention was that all the users of the metadata in the code were removed, so there should be no functional changes. I can remove it though if that would not be considered a NFC.

Is it not possible to just push the cleanupSYCLMetadata pass to run at the end of sycl-post-link?

This is possible, but not all the duplication will be removed during compilation - e.g. the linked device code before sycl-post-link will have duplication in the aspect list but there will be no duplication in the split module files.