libocca / occa

Portable and vendor neutral framework for parallel programming on heterogeneous platforms.
https://libocca.org
MIT License
382 stars 81 forks source link

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

Open tcew opened 5 years ago

tcew commented 5 years ago

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 commented 5 years ago

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 commented 5 years ago

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 commented 5 years ago

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 commented 5 years ago

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

dmed256 commented 5 years ago

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

tcew commented 5 years ago

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 commented 5 years ago

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 commented 5 years ago

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 commented 5 years ago

That sounds reasonable

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

tcew commented 5 years ago

Thanks for contemplating the feature request !