!続・clangのCUDA対応について調査
[[Diary/2011-3-25]]の続き．
/tools/clang/lib/Sema/SemaExpr.cppのSema::ActOnCUDAExecConfigExprの
第3引数MultiExprArg execConfigをみてみる．
MultiExprArgは，tools/clang/include/clang/Sema/Ownership.hで，
 typedef ASTMultiPtr<Expr*> MultiExprArg;
として与えられている型．
  hoge<<<(512), 1>>>(id, Cd);
とかだと，execConfigから，
  execConfig.size()
の値は，2で
  execConfig.get()[0]->dump();
  execConfig.get()[1]->dump();
とかすると，
 (ParenExpr 0x5621eb8 'int'
   (IntegerLiteral 0x5621e90 'int' 512))
 (IntegerLiteral 0x5621ed8 'int' 1)
と得られる．ここまでは，まだ値保存されているね，と一安心．
Sema::ActOnCallExprには第4引数のMultiExprArg argsで引き渡される．
途中で，
  Expr **Args = args.release();
とかってなってる．releaseの定義はtools/clang/include/clang/Sema/Ownership.hに
    PtrTy *release() {
      return Nodes;
    }
とある．Node(Expr型)へのポインタだけを返すようだ．で，
  return BuildResolvedCallExpr(Fn, NDecl, LParenLoc, Args, NumArgs, RParenLoc,
                               ExecConfig);
なので，第4引数で渡される．なので，Sema::BuildResolvedCallExprで，
    Args[0]->dump();
    Args[1]->dump();
としてみてみると，
 (DeclRefExpr 0x42e12b0 'int' lvalue ParmVar 0x42dd110 'id' 'int')
 (DeclRefExpr 0x42e12d8 'float *' lvalue Var 0x42dd2d0 'Cd' 'float *')
...あれ？
  Expr **Args = args.release();
の直後で，Argsをdumpしてみることにする．もちろん，ここは問題ない．
ん？やはり，最後のreturn BuildResolvedCallExpr(...)で返っているわけじゃないのか？
いや，return BuildResolvedCallExpr(...)で返っているのは確かだけど，
  if (Config) {
    TheCall = new (Context) CUDAKernelCallExpr(Context, Fn,
                                               cast<CallExpr>(Config),
                                               Args, NumArgs,
                                               Context.BoolTy,
                                               VK_RValue,
                                               RParenLoc);
  } else {
    TheCall = new (Context) CallExpr(Context, Fn,
                                     Args, NumArgs,
                                     Context.BoolTy,
                                     VK_RValue,
                                     RParenLoc);
  }
ではないみたいだ．フックのいれかたが不適切だったなあ．
というわけで，Configの値をみてみると0なので，else節の方へマッチして，CallExprが呼ばれている．
  if (const FunctionProtoType *Proto = dyn_cast<FunctionProtoType>(FuncT)) {
    if (ConvertArgumentsForCall(TheCall, Fn, FDecl, Proto, Args, NumArgs,
                                RParenLoc))
      return ExprError();
  } else {
    assert(isa<FunctionNoProtoType>(FuncT) && "Unknown FunctionType!");
のthen節にマッチしているようだ．
 (CXXMethodDecl *Method = dyn_cast_or_null<CXXMethodDecl>(FDecl)) = 0
 NDecl = 0x3defb50
 FDecl = 0x3defb50
だけど，結局最後の
  return MaybeBindToTemporary(TheCall);
のとこで返るみたい．
TheCallは，
 (CallExpr 0x4e42f70 '_Bool'
   (ImplicitCastExpr 0x4e42f58 'cudaError_t (*)(dim3, dim3, size_t, cudaStream_t)' <FunctionToPointerDecay>
     (DeclRefExpr 0x4e42f30 'cudaError_t (dim3, dim3, size_t, cudaStream_t)' lvalue Function 0x43f3b50 'cudaConfigureCall' 'cudaError_t (dim3, dim3, size_t, cudaStream_t)'))
   (ParenExpr 0x4e42ee8 'int'
     (IntegerLiteral 0x4e42ec0 'int' 512))
   (IntegerLiteral 0x4e42f08 'int' 1))
なのだけど，返るところの直前では，
 (CallExpr 0x4e42f70 'cudaError_t':'enum cudaError'
   (ImplicitCastExpr 0x4e42f58 'cudaError_t (*)(dim3, dim3, size_t, cudaStream_t)' <FunctionToPointerDecay>
     (DeclRefExpr 0x4e42f30 'cudaError_t (dim3, dim3, size_t, cudaStream_t)' lvalue Function 0x43f3b50 'cudaConfigureCall' 'cudaError_t (dim3, dim3, size_t, cudaStream_t)'))
   (CXXConstructExpr 0x4e430a8 'dim3':'struct dim3''void (const struct dim3 &) throw()' elidable
     (ImplicitCastExpr 0x4e43090 'const struct dim3' <NoOp>
       (ImplicitCastExpr 0x4e43078 'dim3':'struct dim3' <ConstructorConversion>
         (CXXConstructExpr 0x4e43028 'dim3':'struct dim3''void (unsigned int, unsigned int, unsigned int)'
           (ImplicitCastExpr 0x4e42fd0 'unsigned int' <IntegralCast>
             (ParenExpr 0x4e42ee8 'int'
               (IntegerLiteral 0x4e42ec0 'int' 512)))
           (CXXDefaultArgExpr 0x4e42fe8 'unsigned int')
           (CXXDefaultArgExpr 0x4e43008 'unsigned int')))))
   (CXXConstructExpr 0x4e431c0 'dim3':'struct dim3''void (const struct dim3 &) throw()' elidable
     (ImplicitCastExpr 0x4e431a8 'const struct dim3' <NoOp>
       (ImplicitCastExpr 0x4e43190 'dim3':'struct dim3' <ConstructorConversion>
         (CXXConstructExpr 0x4e43140 'dim3':'struct dim3''void (unsigned int, unsigned int, unsigned int)'
           (ImplicitCastExpr 0x4e430e8 'unsigned int' <IntegralCast>
             (IntegerLiteral 0x4e42f08 'int' 1))
           (CXXDefaultArgExpr 0x4e43100 'unsigned int')
           (CXXDefaultArgExpr 0x4e43120 'unsigned int')))))
   (CXXDefaultArgExpr 0x4e43200 'size_t':'unsigned long')
   (CXXDefaultArgExpr 0x4e43220 'cudaStream_t':'struct CUstream_st *'))
となっている．この間に，MaybeBindTemporaryが何ども呼び出されている．
dim3で，512とか，1とか，それぞれちゃんと残っているけど，何の関数なんだ，これ？
あらためて，.llをみてみると，
  %tmp78 = load i32* %id.addr, align 4
  %tmp79 = load float** %Cd, align 8
  call void @_Z6hogeiPf(i32 %tmp78, float* %tmp79)
とかってなっているから，やっぱりどっかで，値なくなってるんだよなあ...
 ExprResult Sema::MaybeBindToTemporary(Expr *E);
はinclude/clang/Sema/Sema.hに定義があって，lib/Sema/SemaExprCXX.cppに実装がある．
  const RecordType *RT = E->getType()->getAs<RecordType>();
  if (!RT)
    return Owned(E);
のところでマッチして返るみたい．たしかにRecordTypeではないよなあ．
Ownedは，
  ExprResult Owned(Expr* E) { return E; }
はinclude/clang/Sema/Sema.hに．ちなみにExprResultの定義は，include/clang/Sema/Ownership.hに．
 typedef ActionResult<Expr*> ExprResult;
なるほど．cudaConfigureCallっていう関数呼び出しになるのか...
いづれにしても，でてこないけど．