Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Feature request: add option to stop parser for replacing defined variables with their values #180

Open
tcew opened this issue Jul 29, 2018 · 10 comments
Labels
feature Use this label to request a new feature! parser

Comments

@tcew
Copy link
Collaborator

tcew commented Jul 29, 2018

Currently if p_N is defined to 5 then when a kernel is parsed and output all instances of p_N are replaced by 5.

It would be useful to add a parser option so that the #defines are included in the generated source code, and the compilers variables like p_N are left in the generated source code.

This will make the generated code a lot easier to read.

@dmed256
Copy link
Member

dmed256 commented Jul 29, 2018

Unfortunately all defines need to be expanded to apply code transformations
#define values can be found in two places

raw_source.cpp

> head ~/.occa/cache/c8141715ac4e4272/raw_source.cpp 
#define  block 256

/* The MIT License (MIT)
 *

build.json

> cat ~/.occa/cache/c8141715ac4e4272/build.json | jq .kernel.props.defines
{
  "block": 256
}

@tcew
Copy link
Collaborator Author

tcew commented Jul 29, 2018

I agree that it is useful to do this to extract actual loop bounds and I imagine the parser is designed to do this. However, it seems unlikely that it is impossible to preserve the definitions in the code. Keeping track of how the values are computed would obviate the need to print their numerical values.

@dmed256
Copy link
Member

dmed256 commented Jul 29, 2018

It is harder than it seems to maintain the code and transform it
Defines can be anything, even partial code

#define foo 3 +

for (int i = 0; i < foo 5; ++i) {}

Keeping the #define lines can also cause issues since the compiler will still run the preprocessor

@tcew
Copy link
Collaborator Author

tcew commented Jul 29, 2018

... and if the Define is just a numerical constant... ?

@dmed256
Copy link
Member

dmed256 commented Jul 29, 2018

... and why do we need extra logic for that ....?

@tcew
Copy link
Collaborator Author

tcew commented Jul 29, 2018

This is an example of the output kernel

extern "C" __global__ void _occa_ellipticPreconCoarsenHex3D_0(const int Nelements,
                                                              const double * __restrict__ R,
                                                              const double * __restrict__ qf,
                                                              double * __restrict__ qc) {
  {
    int e = 0 + blockIdx.x;
    __shared__ double s_qfff[8][8][8];
    __shared__ double s_qcff[2][8][8];
    __shared__ double s_qccf[2][2][8];
    __shared__ double s_qccc[2][2][2];
    __shared__ double s_R[2][8];
    {
      int k = 0 + threadIdx.z;
      {
        int j = 0 + threadIdx.y;
        {
          int i = 0 + threadIdx.x;
          const int id = i + j * 8 + k * 8 * 8 + e * 512;
          s_qfff[k][j][i] = qf[id];
          if ((k == 0) && (j < 2)) {
            s_R[j][i] = R[j * 8 + i];
          }
        }
      }
    }

The kernels are losing readability on translation. All these numbers were Defines.

@dmed256
Copy link
Member

dmed256 commented Jul 30, 2018

Taking a step back, the reason for the feature request was to

make the generated code a lot easier to read

The generated code is meant for the compiler, not for users (similar to preprocessed outputs by the compiler). If there is a parsing issue, specially since the parser is not yet mature, it could help to look at. However, these issues should be sorted out over time.

If the purpose is to make it easier to debug, I agree we should add #line to match with the original source code

@tcew
Copy link
Collaborator Author

tcew commented Jul 30, 2018

I agree that for debugging this would be a true bonus - nice idea !

I should have been clearer about the reasoning behind this request: one important use case I envision for the new code generation tools in OCCA 1.0 is as a translator from OKL to native threading language for non-OCCA applications. In that case it will be important to preserve readability as the generated CUDA/OpenCL/HIP/OpenMP code will be separated from the OKL code.

@dmed256
Copy link
Member

dmed256 commented Jul 30, 2018

That sounds reasonable

There are a few features that will need to be added first

  • Create global variables from numerical defines passed as kernel properties rather than #define
    • ⚠️ Being able to do it for all #defines is too difficult since we would need to track where they are defined and which scopes they touch after transforms
  • [New Feature] During parser::getExpression, substitute constexpr with their values
    • Initially only const variables
    • Expand to constexpr functions
  • Add back #define lines to the top like usual

@dmed256 dmed256 added feature Use this label to request a new feature! parser and removed Question labels Jul 30, 2018
@tcew
Copy link
Collaborator Author

tcew commented Jul 30, 2018

Thanks for contemplating the feature request !

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
feature Use this label to request a new feature! parser
Projects
None yet
Development

No branches or pull requests

2 participants