Broken behavior with target attribute and enqueued blocks

[arsenm]

It's not clear to me what the behavior of the target attribute is supposed to be with enqueued blocks, but what it does now is broken and inconsistent.

__attribute__((target("s-memtime-inst")))
kernel void test_target_features_kernel(global int *i) {
  queue_t default_queue;
  unsigned flags = 0;
  ndrange_t ndrange;

  __builtin_amdgcn_s_memtime();

  enqueue_kernel(default_queue, flags, ndrange,
                 ^(void) {
                   __builtin_amdgcn_s_memtime();
                 });
}

This program is currently accepted. If I remove the target attribute from the kernel, the builtin use in the kernel code is correctly rejected. The builtin use in the block is never diagnosed. If you attempt to apply the target attribute to the block, it's disallowed. So either one of these must be true:

1. The target attribute should propagate to the block, in which case the implied target-features entries should propagate to the block in the IR
2. If the target attribute is not implied on the block, the target attribute should probably be allowed on the block itself.

The diagnostics for the builtin in the block are missing in either case

[llvmbot]

@llvm/issue-subscribers-opencl

[llvmbot]

@llvm/issue-subscribers-clang-codegen