add enzyme_noatomic attribute

[minansys]

Preserve and honor enzyme_noatomic on CUDA functions/args

Add support for the "enzyme_noatomic" annotation in PreserveNVVM and
teach the reverse-mode atomic path to disable atomics when either the
function or a base pointer argument carries that attribute. This keeps
CUDA annotations round-tripped into LLVM attributes and applies them
consistently during gradient generation.

Usage: mark a kernel or helper with attribute((annotate("enzyme_noatomic")))
(or attach "enzyme_noatomic" to a specific pointer argument), run
preserve-nvvm before Enzyme, and Enzyme will avoid atomicrmw in the
generated reverse kernel.

// Minimal CUDA kernel to exercise enzyme_noatomic via annotations.
#include <cuda_runtime.h>
#define ENZYME_NOATOMIC __attribute__((annotate("enzyme_noatomic")))
extern "C" {
__device__ int enzyme_dup;
__device__ void __enzyme_autodiff(void *, ...);
}
extern "C" __device__ ENZYME_NOATOMIC void my_kernel_body(float *out,
                                                          const float *in) {
  uint tid = blockIdx.x * blockDim.x + threadIdx.x;
  out[0] += in[tid];
}
extern "C" __device__ void test_derivative(float *out, float *dout,
                                           const float *in, float *din) {
  __enzyme_autodiff((void *)my_kernel_body, enzyme_dup, out, dout, enzyme_dup,
                    in, din);
}

the result with and without the ENZYME_NOATOMIC

[Copilot]

Pull request overview

Adds a new Enzyme annotation/attribute plumbing path intended to influence reverse-mode accumulation on GPU targets by disabling atomic updates when a marker attribute is present, and preserves that marker through the NVVM preservation pass.

Changes:

• Teach PreserveNVVM to convert the llvm.global.annotations string "enzyme_elementwise_read" into an LLVM function attribute.
• Add a Clang plugin parsed attribute enzyme_elementwise_read that lowers to AnnotateAttr("enzyme_elementwise_read").
• Disable reverse-mode atomic accumulation in DiffeGradientUtils::addToInvertedPtrDiffe when the function (and nominally a base-pointer argument) carries "enzyme_elementwise_read".

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 4 comments.

File	Description
enzyme/Enzyme/PreserveNVVM.cpp	Preserves a new annotation string by translating it into an LLVM function attribute.
enzyme/Enzyme/DiffeGradientUtils.cpp	Adds an attribute check that can disable atomic derivative accumulation based on context.
enzyme/Enzyme/Clang/EnzymeClang.cpp	Adds Clang-side parsing for enzyme_elementwise_read and lowers it to AnnotateAttr.

[Copilot]

This change alters reverse-mode accumulation behavior by disabling atomics when the new attribute is present, but there’s no lit/IR test exercising the new attribute path (e.g., ensuring PreserveNVVM round-trips the annotation into an LLVM fn attribute and that the generated reverse kernel avoids atomicrmw for that accumulation). Please add a regression test covering both the attribute preservation and the expected removal of atomicrmw in the differentiated output.

[minansys]

@copilot apply changes based on this feedback

[minansys]

@copilot fix this clang format issue "--- enzyme/Enzyme/PreserveNVVM.cpp (original)
+++ enzyme/Enzyme/PreserveNVVM.cpp (reformatted)
@@ -388,10 +388,9 @@
}

         if (AS == "enzyme_elementwise_read" && Func) {

•          Func->addAttribute(
  

•              AttributeList::FunctionIndex,
  

•              Attribute::get(Func->getContext(),
  

•                             "enzyme_elementwise_read"));
  

•          Func->addAttribute(AttributeList::FunctionIndex,
  

•                             Attribute::get(Func->getContext(),
  

•                                            "enzyme_elementwise_read"));
           changed = true;
           replacements.push_back(Constant::getNullValue(CAOp->getType()));
           continue;"