[HIPIFY] Kernel launch inside a macro

[incardon]

Having a kernel launch inside a macro like this

#define CUDA_LAUNCH(cuda_call,wthr,thr, ...) \
		cuda_call<<<wthr,thr>>>(__VA_ARGS__);

and using it with

CUDA_LAUNCH((count_nn_cells),itgg,vs.toKernel(),cells_nn.toKernel(),cells_nn_test.toKernel());

produce a crash. The main problem come from HipifyAction.cpp

size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart);

It seem that launchEnd point to a location before launchStart

printing the location pointed by the two positions /tmp/CellList_gpu_test.cu-b82dd5.hip:555:2 <----- end /tmp/CellList_gpu_test.cu-b82dd5.hip:555:14 <----- start

producing a length negative and subsequent crash.

I use this fact to skip the substitution. ( That is also what I want, because is the macro to be changed and not where the macro is called )

  long int length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart);

  if (length > 0)
  {
        ct::Replacement Rep(*SM, launchStart, length, OS.str());
        clang::FullSourceLoc fullSL(launchStart, *SM);
        insertReplacement(Rep, fullSL);
        hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_EXECUTION, ApiTypes::API_RUNTIME};
        Statistics::current().incrementCounter(counter, refName.str());
  }

I report it in case there is a better way to fix the problem

[emankov]

Could you please provide the whole test case?

[incardon]

I lanuch with

../HIP/hipify-clang/dist/bin/hipify-clang -I/usr/local/cuda/include -I. -I./config -I/where/you/have/BOOST_1.68_or_higher/include -DNDEBUG NN/CellList/CellList_gpu_test.cu -D__NVCC__ -o CellList_gpu_test.cpp -D__clang__ -- --std=c++11 -D__CUDACC_VER_MAJOR__=10 -D__CUDACC_VER_MINOR__=1 -D__CUDACC_VER_MINOR__ -D__CUDACC_VER_BUILD__=243

You have to set boost. I know is not the smallest test case, but if you are able to reproduce from here, you save me time in searching the smallest case to reproduce. If you think it is too big I will create the smallest one.

sources_bug.zip

[emankov]

Hi @incardon,

I've just fixed it finally, but not the way you thought. Please, have a look.

That is also what I want, because is the macro to be changed and not where the macro is called

You right, but this particular macro is really something: __VA_ARGS__ and function style - things, expanding which I have not overcome yet. Please, have a look at the test: https://github.com/ROCm-Developer-Tools/HIP/blob/master/tests/hipify-clang/unit_tests/kernel_launch/kernel_launch_01.cu as an interim working solution. If the solution suits you, please close this Issue, and I will file a new one for replacement in macro.

[emankov]

@incardon, just a gentle ping.

[incardon]

Thank you very much for the fix. I will test it as soon as possible

[emankov]

@incardon, have you had a chance to test on your side?

[incardon]

Yes. Thanks. It does the Job perfectly

[emankov]

But it is not how you wanted initially.

[incardon]

If you mean it skip the substitution. It does. If you refer to the fact that does not change the macro it does not, but at least in my case is marginal problem. It is just one place where I have to change the code manually. That macro is just used to do error checking in debug mode

[emankov]

I meant that the proper solution is to transform the macro declaration, not its usage transforming. ROCm-Developer-Tools/HIP#1409 is a compromise. I'll open a new ticket not to forget implementing macro declaration transformation.