[PATCH] D25343: [OpenCL] Mark group functions as convergent in opencl-c.h

[Ettore Speziale via cfe-commits]

Hello guys,

>> Should we deprecate noduplicate then as convergent should cover both use cases for OpenCL I believe? As far as I understand noduplicate was added specifically for SPMD use cases...
> 
> noduplicate has different semantics than convergent. Although it is proposed for SPMD originally, it could be used by non-SPMD programs to forbid duplicate of functions. There may be applications using this attribute.
> 
> I would suggest to leave this question for future. Probably ask llvm-dev first since the attribute is also in LLVM.

I just want to clarify why I withdraw the convergent patch I initially submitted some time ago. It has a problem when dealing with multiple modules. Consider the following example:

int foo(void) __attribute__((pure));

int bar(int x) {
  int y = foo();
  if (x)
    return y;
  return 0;
}   

I’ve just marked foo with the pure attribute to mark the function readonly in LLVM IR. Given that IR, the IR sinking pass pushes the foo call site into the then branch:

; Function Attrs: nounwind readonly ssp uwtable
define i32 @bar(i32) #0 {
  %2 = icmp eq i32 %0, 0
  br i1 %2, label %5, label %3

; <label>:3                                       ; preds = %1
  %4 = tail call i32 @foo() #2
  br label %5

; <label>:5                                       ; preds = %1, %3
  %6 = phi i32 [ %4, %3 ], [ 0, %1 ]
  ret i32 %6
}

; Function Attrs: nounwind readonly
declare i32 @foo() #1

This is kind of dangerous, as we do not know what is inside foo — i.e. it might contains a convergent call.

If I understand correctly, the CUDA guys solved the problem in two steps. At CodeGen time all the device function calls are marked convergent:

  if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
    // Conservatively, mark all functions and calls in CUDA as convergent
    // (meaning, they may call an intrinsically convergent op, such as
    // __syncthreads(), and so can't have certain optimizations applied around
    // them).  LLVM will remove this attribute where it safely can.
    FuncAttrs.addAttribute(llvm::Attribute::Convergent);

Then LLVM function attribute pass — lib/Transforms/IPO/FunctionAttrs.cpp — remove the unnecessary convergent attributes starting from the leaf nodes  — i.e. external calls.

Provide that intrinsics are correctly marked convergent only when needed, that allow to get rid of the unnecessary convergent attributes.

Since you are introducing an explicit convergent attribute it seems that OpenCL is requiring the developers to explicitly mark the functions that might contain convergent function calls with the convergent attribute. Am I understand correctly?

Thanks

--------------------------------------------------
Ettore Speziale — Compiler Engineer
speziale.ettore at gmail.com
espeziale at apple.com
--------------------------------------------------