[OpenCL] Always add nounwind attribute for OpenCL

[svenvh]

Summary

Neither OpenCL nor C++ for OpenCL support exceptions, so add the
nounwind attribute unconditionally for those languages.

Unblocks D138958.

[lebedev.ri]

Thanks!

[bader]

Should we generalize and rename clang/test/CodeGenOpenCL/convergent.cl to validate function attributes other than convergent? It's not obvious that presence of nounwind attribute is validated by clang/test/CodeGenOpenCL/convergent.cl.

[svenvh]

In D142033#4062671, @bader wrote:

Should we generalize and rename clang/test/CodeGenOpenCL/convergent.cl to validate function attributes other than convergent? It's not obvious that presence of nounwind attribute is validated by clang/test/CodeGenOpenCL/convergent.cl.

I think the main goal of clang/test/CodeGenOpenCL/convergent.cl remains testing convergent, so I'd rather not generalize this particular test. nounwind is somewhat coincidentally tested by various tests now, but we could add a separate test for generic attributes such as nounwind if you think it's worth doing so.

clang/lib/CodeGen/CGCall.cpp

if (getLangOpts().assumeFunctionsAreConvergent()) {
     // Conservatively, mark all functions and calls in CUDA and OpenCL as
     // convergent (meaning, they may call an intrinsically convergent op, such
     // as __syncthreads() / barrier(), and so can't have certain optimizations
     // applied around them).  LLVM will remove this attribute where it safely
     // can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
   }
 
-  // TODO: NoUnwind attribute should be added for other GPU modes OpenCL, HIP,
+  // TODO: NoUnwind attribute should be added for other GPU modes HIP,
   // SYCL, OpenMP offload. AFAIK, none of them support exceptions in device
   // code.
-  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
-    // Exceptions aren't supported in CUDA device code.
+  if ((getLangOpts().CUDA && getLangOpts().CUDAIsDevice) ||
+      getLangOpts().OpenCL) {
     FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
   }
 
   for (StringRef Attr : CodeGenOpts.DefaultFunctionAttrs) {
     StringRef Var, Value;
     std::tie(Var, Value) = Attr.split('=');
     FuncAttrs.addAttribute(Var, Value);
   }

clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl

 // CHECK-NEXT:    ret void
 //
 //.
 // CHECK: attributes #0 = { convergent noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
 // CHECK: attributes #1 = { convergent noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,256" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" "uniform-work-group-size"="false" }
 // CHECK: attributes #2 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
 // CHECK: attributes #3 = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "target-features"="+16-bit-insts,+ci-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
 // CHECK: attributes #4 = { nounwind "enqueued-block" }
-// CHECK: attributes #5 = { convergent }
+// CHECK: attributes #5 = { convergent nounwind }
 //.
 // CHECK: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
 // CHECK: !1 = !{i32 1, !"wchar_size", i32 4}
 // CHECK: !2 = !{i32 2, i32 0}
 // CHECK: !3 = !{i32 1, i32 0, i32 1, i32 0}
 // CHECK: !4 = !{!"none", !"none", !"none", !"none"}
 // CHECK: !5 = !{!"char*", !"char", !"long*", !"long"}
 // CHECK: !6 = !{!"", !"", !"", !""}

clang/test/CodeGenOpenCL/convergent.cl

 }
 
 // CHECK: attributes #0 = { nofree noinline norecurse nounwind "
 // CHECK: attributes #1 = { {{[^}]*}}convergent{{[^}]*}} }
 // CHECK: attributes #2 = { {{[^}]*}}convergent{{[^}]*}} }
 // CHECK: attributes #3 = { {{[^}]*}}convergent noduplicate{{[^}]*}} }
 // CHECK: attributes #4 = { {{[^}]*}}convergent{{[^}]*}} }
 // CHECK: attributes #5 = { {{[^}]*}}convergent{{[^}]*}} }
-// CHECK: attributes #6 = { {{[^}]*}}convergent noduplicate{{[^}]*}} }
+// CHECK: attributes #6 = { {{[^}]*}}nounwind{{[^}]*}} }
+// CHECK: attributes #7 = { {{[^}]*}}convergent noduplicate nounwind{{[^}]*}} }