[AMDGPU] fixes for lds f32 builtins

[dfukalov]

Summary

1. added restrictions to memory scope, order and volatile parameters
2. added custom processing for these builtins - currently is not used code, needed to switch off GCCBuiltin link to the builtins (ongoing change to llvm tree)
3. builtins renamed as requested

[b-sumner]

lib/CodeGen/CGBuiltin.cpp

10094

Can the pointer argument address space be checked here?

[dfukalov]

diff updated as requested by reviewer

[dfukalov]

lib/CodeGen/CGBuiltin.cpp

10094

added code to check pointer and flaot parameters
though this part is not active at the moment since the builtins are not processed as custom.
going to switch to custom path by removing attribute of the intrinsics in llvm trunk in next step

[dfukalov]

ping...

[arsenm]

Can’t you just change the description to be the LangAS value? I also thought these happened to be the same already

[dfukalov]

In D43281#1018657, @arsenm wrote:

Can’t you just change the description to be the LangAS value? I also thought these happened to be the same already

Am I right that you mean to change the semantic of the addrspace number in a description string for all targets?

At the moment it's finally checked in ASTContext::getAddrSpaceQualType that LangAS is equal addrspace returned by QualType.
And for addrspace "2" specified in description, this QualType is "attribute((address_space(3))) float". And returns "11" since it's target addrspace, defined in LangAS after FirstTargetAddressSpace.
So ASTContext::getAddrSpaceQualType goes through check CanT.getAddressSpace() == AddressSpace and then hits assertions that the parameter already has addrspace specified.

[arsenm]

I meant we could just change the value to whatever it happens to be for AMDGPU. That this is the language address space is surprising to me though, so maybe that should change

[dfukalov]

The problem is that if set addrspace "2" in description string, CanT.getAddressSpace() returns target addrspace value "11" (shifted in the enum) and compares it with input LangAS addrspace ("2", "opencl_local" in our case).
So I cannot set a number a description string that will be equal to LangAS addrspace "opencl_local".

Moreover, this change is preparation to move to custom processing of these builtins. Then I'm going to remove link (GCCBuiltin in IntrinsicsAMDGPU.td) from the llvm intrinsics definitions. Then I'll be able to switch on custom processing in cfe.

[dfukalov]

addrspace specifications are kept in descriptions strings

[dfukalov]

ping...

[arsenm]

In D43281#1023962, @dfukalov wrote:

The problem is that if set addrspace "2" in description string, CanT.getAddressSpace() returns target addrspace value "11" (shifted in the enum) and compares it with input LangAS addrspace ("2", "opencl_local" in our case).
So I cannot set a number a description string that will be equal to LangAS addrspace "opencl_local".

Moreover, this change is preparation to move to custom processing of these builtins. Then I'm going to remove link (GCCBuiltin in IntrinsicsAMDGPU.td) from the llvm intrinsics definitions. Then I'll be able to switch on custom processing in cfe.

My real question was what happens if you put 11 in the description string?

[dfukalov]

My real question was what happens if you put 11 in the description string?

in this case CanT.getAddressSpace() returns target addrspace value "20" (also shifted in the enum by 9==LangAS::FirstTargetAddressSpace)

So again ASTContext::getAddrSpaceQualType decieds that these addrspaces are not equal and then asserts
a
The main reason of these change is to allow usage of the builtins with __local described pointers. Such pointers have LangAS AddressSpace in the ASTContext::getAddrSpaceQualType() equal to "2".
But we have no way to obtain this value from CanT.getAddressSpace() since it returns target addrspace and it is >= 10.

So I'm going to switch to custom processing of these builtins and the patch is preparation to be able to remove their links in llvm repo.

[dfukalov]

ping...

[dfukalov]

ping...

[yaxunl]

LGTM. Thanks!

[arsenm]

I'm looking at how the address space mapping works for builtins, and I think what's there is just uselessly broken and needs to be fixed. It seems to be operating under the assumption that the address spaces the target defines are totally disjoint from the language address spaces

lib/CodeGen/CGBuiltin.cpp

10100–10104

Condition backwards

[arsenm]

I think the intent of the current code is for the address space to correspond to a "target address space" as if the user code used attribute((address_space(n))) to specify a pointer value. This is confusingly named, and different from the target address space selected for a LangAS. I think we need to add some mechanism for specifying the builtin is a LangAS ID. Since Ideally this would also work for multiple languages (e.g. cuda_constant or opencl_constant for the same builtin) I think there needs to be some callback triggered for the address space value. This possibly needs to be distinct from the current pointer descriptor to avoid breaking the possibility of user defined address spaces. There aren't really any other users of builtins with address spaces. NVPTX has some, but the tests seem to not actually try to use the declared address space and pass generic pointers to them.

include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_fcmp, "LUiddIi", "nc")
 BUILTIN(__builtin_amdgcn_fcmpf, "LUiffIi", "nc")
 BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi", "nc")
 BUILTIN(__builtin_amdgcn_ds_permute, "iii", "nc")
 BUILTIN(__builtin_amdgcn_ds_bpermute, "iii", "nc")
 BUILTIN(__builtin_amdgcn_readfirstlane, "ii", "nc")
 BUILTIN(__builtin_amdgcn_readlane, "iii", "nc")
 BUILTIN(__builtin_amdgcn_fmed3f, "ffff", "nc")
-BUILTIN(__builtin_amdgcn_ds_fadd, "ff*3fiib", "n")
-BUILTIN(__builtin_amdgcn_ds_fmin, "ff*3fiib", "n")
-BUILTIN(__builtin_amdgcn_ds_fmax, "ff*3fiib", "n")
+BUILTIN(__builtin_amdgcn_ds_faddf, "ff*fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fminf, "ff*fIiIiIb", "n")
+BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*fIiIiIb", "n")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh", "nc", "16-bit-insts")
 TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh", "nc", "16-bit-insts")

lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
     StringRef RegName = BuiltinID == AMDGPU::BI__builtin_amdgcn_read_exec_lo ?
       "exec_lo" : "exec_hi";
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int32Ty, Int32Ty, true, RegName));
     CI->setConvergent();
     return CI;
   }
+  case AMDGPU::BI__builtin_amdgcn_ds_faddf:
+  case AMDGPU::BI__builtin_amdgcn_ds_fminf:
+  case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: {
+    llvm::SmallVector<llvm::Value *, 5> Args;
+    for (unsigned I = 0; I != 5; ++I)
+      Args.push_back(EmitScalarExpr(E->getArg(I)));
+    const llvm::Type *PtrTy = Args[0]->getType();
+    // check pointer parameter
+    if (!PtrTy->isPointerTy() ||
+        E->getArg(0)
+                ->getType()
+                ->getPointeeType()
+                .getQualifiers()
+                .getAddressSpace() != LangAS::opencl_local ||
+        !PtrTy->getPointerElementType()->isFloatTy()) {
+       CGM.Error(E->getArg(0)->getLocStart(),
+                "parameter should have type \"local float*\"");
+      return nullptr;
+    }
+    // check float parameter
+    if (!Args[1]->getType()->isFloatTy()) {
+      CGM.Error(E->getArg(1)->getLocStart(),
+                "parameter should have type \"float\"");
+      return nullptr;
+    }
+
+    Intrinsic::ID ID;
+    switch (BuiltinID) {
+    case AMDGPU::BI__builtin_amdgcn_ds_faddf:
+      ID = Intrinsic::amdgcn_ds_fadd;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_ds_fminf:
+      ID = Intrinsic::amdgcn_ds_fmin;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
+      ID = Intrinsic::amdgcn_ds_fmax;
+      break;
+    default:
+      llvm_unreachable("Unknown BuiltinID");
+    }
+    Value *F = CGM.getIntrinsic(ID);
+    return Builder.CreateCall(F, Args);
+  }
 
   // amdgcn workitem
   case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0, 1024);
   case AMDGPU::BI__builtin_amdgcn_workitem_id_y:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_y, 0, 1024);
   case AMDGPU::BI__builtin_amdgcn_workitem_id_z:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 1024);

test/CodeGenOpenCL/builtins-amdgcn-vi.cl

 // CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1 false)
 void test_mov_dpp(global int* out, int src)
 {
   *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
 }
 
 // CHECK-LABEL: @test_ds_fadd
 // CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fadd(__attribute__((address_space(3))) float *out, float src) {
-  *out = __builtin_amdgcn_ds_fadd(out, src, 0, 0, false);
+void test_ds_faddf(local float *out, float src) {
+  *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false);
 }
 
 // CHECK-LABEL: @test_ds_fmin
 // CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fmin(__attribute__((address_space(3))) float *out, float src) {
-  *out = __builtin_amdgcn_ds_fmin(out, src, 0, 0, false);
+void test_ds_fminf(local float *out, float src) {
+  *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
 }
 
 // CHECK-LABEL: @test_ds_fmax
 // CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
-void test_ds_fmax(__attribute__((address_space(3))) float *out, float src) {
-  *out = __builtin_amdgcn_ds_fmax(out, src, 0, 0, false);
+void test_ds_fmaxf(local float *out, float src) {
+  *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
 }

test/SemaOpenCL/builtins-amdgcn-error.cl

 void test_mov_dpp2(global int* out, int a, int b, int c, int d, bool e)
 {
   *out = __builtin_amdgcn_mov_dpp(a, b, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
   *out = __builtin_amdgcn_mov_dpp(a, 0, c, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
   *out = __builtin_amdgcn_mov_dpp(a, 0, 0, d, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
   *out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
 }
 
+void test_ds_faddf(local float *out, float src, int a) {
+  *out = __builtin_amdgcn_ds_faddf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_faddf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
+}
+
+void test_ds_fminf(local float *out, float src, int a) {
+  *out = __builtin_amdgcn_ds_fminf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fminf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
+}
+
+void test_ds_fmaxf(local float *out, float src, int a) {
+  *out = __builtin_amdgcn_ds_fmaxf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
+  *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
+}