Recompute the values of CUDA built-in index functions

[ovdiiuv]

Before this PR, we used to save the values like threadIdx.x, see exapmle.

__global__ void add_kernel_3(int *out, int *in) {
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  out[index] += in[index];
}

void add_kernel_3_grad(int *out, int *in, int *_d_out, int *_d_in) {
    unsigned int _t1 = blockIdx.x;
    unsigned int _t0 = blockDim.x;
    int _d_index = 0;
    int index0 = threadIdx.x + _t1 * _t0;
    int _t2 = out[index0];
    out[index0] += in[index0];
    {
        out[index0] = _t2;
        int _r_d0 = _d_out[index0];
        atomicAdd(&_d_in[index0], _r_d0);
    }
}

Here we create _t0 to store the blockDim.x, but clearly the function doesn't have any side effects, hence should be recomputed.

[github-actions]

clang-tidy made some suggestions

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

📢 Thoughts on this report? Let us know!

[github-actions]

clang-tidy review says "All clean, LGTM! 👍"

[kchristin22]

Instead of checking if this is a CUDA Builtin, since effectively these are const variables, can we extend this functionality for all const expressions?

[ovdiiuv]

Yes, you are right. I took this approach because extracting CallExpr from lets say threadIdx.x is rather ugly, so I guess it deserves to be in a separate function.

[kchristin22]

Yes, I agree with this being a separate function, but instead of checking the names we can check if it's a const type. In the CUDA case, you can make it to check directly the second subE (OpaqueValueExpr) whether it's a const type. I guess this could work and you wouldn't have to dive that deep as in the case of the CallExpr. But it would be nice if we had a more generic function that checks if the expression is of const type, and one if case there would include this PseudoObjectExpr check-path.

[ovdiiuv]

Yes, I agree.

[ovdiiuv]

Addressed. I think we would need to look into a solution for a more generic case, which is out of the scope of this PR.

[kchristin22]

Maybe something like this?

if (const auto *pseudoE = llvm::dyn_cast<PseudoObjectExpr>(B)) {
  if (const auto *opaqueE =
          llvm::dyn_cast<OpaqueValueExpr>(pseudoE->getSemanticExpr(0))) {
      const Expr *innerE = opaqueE->getSourceExpr()->IgnoreImplicit();
      QualType innerT = innerE->getType();
      if (innerT.isConstQualified())
          return true;
  }
 }

[ovdiiuv]

Looks so much better! Thanks:)

[vgvassilev]

This can be a file static function.

[github-actions]

clang-tidy review says "All clean, LGTM! 👍"

[github-actions]

clang-tidy review says "All clean, LGTM! 👍"

[github-actions]

clang-tidy review says "All clean, LGTM! 👍"

[github-actions]

clang-tidy review says "All clean, LGTM! 👍"

[github-actions]

clang-tidy review says "All clean, LGTM! 👍"