kokkos/llvm-project

The LLVM Project is a collection of modular and reusable compiler and toolchain technologies. Note: the repository does not accept github pull requests at this moment. Please submit your patches at http://reviews.llvm.org.

Kokkos: Core Libraries

Kokkos Core implements a programming model in C++ for writing performance portable applications targeting all major HPC platforms. For that purpose it provides abstractions for both parallel execution of code and data management. Kokkos is designed to target complex node architectures with N-level memory hierarchies and multiple types of execution resources. It currently can use CUDA, HIP, SYCL, HPX, OpenMP and C++ threads as backend programming models with several other backends in development.

Kokkos Core is part of the Kokkos C++ Performance Portability Programming Ecosystem.

Kokkos is a Linux Foundation project.

Learning about Kokkos

To start learning about Kokkos:

• Kokkos Lectures: they contain a mix of lecture videos and hands-on exercises covering all the important capabilities.

• Programming guide: contains in "narrative" form a technical description of the programming model, machine model, and the main building blocks like the Views and parallel dispatch.

• API reference: organized by category, i.e., core, algorithms and containers or, if you prefer, in alphabetical order.

• Use cases and Examples: a serie of examples ranging from how to use Kokkos with MPI to Fortran interoperability.

Obtaining Kokkos

The latest release of Kokkos can be obtained from the GitHub releases page.

The current release is 4.3.01.

curl -OJ -L https://github.com/kokkos/kokkos/archive/refs/tags/4.3.01.tar.gz
# Or with wget
wget https://github.com/kokkos/kokkos/archive/refs/tags/4.3.01.tar.gz

To clone the latest development version of Kokkos from GitHub:

git clone -b develop  https://github.com/kokkos/kokkos.git

Building Kokkos

To build Kokkos, you will need to have a C++ compiler that supports C++17 or later. All requirements including minimum and primary tested compiler versions can be found here.

Building and installation instructions are described here.

You can also install Kokkos using Spack: spack install kokkos. Available configuration options can be displayed using spack info kokkos.

For the complete documentation: kokkos.org/kokkos-core-wiki/

Support

For questions find us on Slack: https://kokkosteam.slack.com or open a GitHub issue.

For non-public questions send an email to: crtrott(at)sandia.gov

Contributing

Please see this page for details on how to contribute.

Citing Kokkos

Please see the following page.

License

Under the terms of Contract DE-NA0003525 with NTESS, the U.S. Government retains certain rights in this software.

The full license statement used in all headers is available here or here.

llvm-project's Issues

Figure out if it is even possible to do clang-tidy plugins

This doesn't currently exist, but it would be really nice and possibly very upstreamable.

warnings in uninstantiated templates

see kokkos/kokkos#2579 (comment)

@vbrunini

I haven't forgotten about you.

Add RAJA check for dereferencing struct pointers in a RAJA loop

This would be a check to see if a RAJA lambda has captured a pointer to a struct and then dereferences it within the kernel. This is the same basic problem that we have when implicitly dereferencing members off of the "this" pointer.

So something like:

struct mystruct* mystruct_ptr = ... ;

struct mystruct mystruct_captured = ...;

RAJA::for_all (RAJA::RangeSegment(0, n, [=](int i) {
...
data[i] += mystruct_ptr->myfield; // catch this dereference

data[i] += mystruct_captured.myfield; // this shouldn't be caught since it's not a pointer dereference
...
} );

__clang_analyzer__ for clangd

Enable the __clang_analyzer__ definition for clangd. I'm not sure this is really the correct thing to do, but as we have things set up now it is the only way to get ensure-kokkos-function to work correctly with clangd.

This is because Kokkos only annotates KOKKOS_.*FUNCITON when the __clang_analyzer__ definition is detected.

I think a longer term solution is to add something like __kokkos_analyzer__ to clangd/clang-tidy so that we know we are in a context that Kokkos definitions should be modified.

@DavidPoliakoff thoughts

Warn if kernel/fence does not have a name

This one shouldn't be too bad, something to the effect of

match callExpr(callee(functionDecl(namedDecl(matchesName("name_we_care_about")))),unless(hasAnyArgument(isStringish)))

Or I guess you could match the types of the parameters

match callExpr(callee(functionDecl(namedDecl(matchesName("name_we_care_about")),unless(hasAnyParameter(isStringish)))))

Guarantee that within Kokkos we name all of our kernels and fences

Thyra specific checks

TODO. @bartlettroscoe sorry for the long delay in getting back to this.

Modify kokkos-implicit-this-capture to warn at the variable usage

I think this warning would be more clear if the location given by the warning was the location of the used variable instead of the offending lambda.

1. Make the warning the actual variable
2. Make the note the location of the lambda.

Change name of default branch from master to something else

Waiting to see what upstream does about this, but wanted to make a note to remind me to do this eventually.

parallel_reduce require reference parameter

We need to make sure in the following case that the out parameter is actually a reference.

Kokkos::parallel_reduce(n,
  KOKKOS_LAMBDA (int i, double out) {
    out += i;
  },
  result
);

check functor arguments to parallel_X functions

Need to detect that the functor passed into to parallel_X function is valid, i.e. the operator called either has KOKKOS_INLINE_FUNCTION or is a KOKKOS_LAMBDA

kokkos-implicit-this-capture warns for explicit *this capture

Using KOKKOS_CLASS_LAMBDA with Kokkos' clang-tidy gives

/var/jenkins/workspace/Kokkos/containers/unit_tests/TestErrorReporter.hpp:198:9: error: Lambda passed to parallel_for implicitly captures this. [kokkos-implicit-this-capture,-warnings-as-errors]
        KOKKOS_CLASS_LAMBDA(const int work_idx) {
        ^
/var/jenkins/workspace/Kokkos/core/src/setup/Kokkos_Setup_Cuda.hpp:87:29: note: expanded from macro 'KOKKOS_CLASS_LAMBDA'
#define KOKKOS_CLASS_LAMBDA [ =, *this ] __host__ __device__
                            ^
/var/jenkins/workspace/Kokkos/containers/unit_tests/TestErrorReporter.hpp:198:9: note: the captured variable is used here.
        KOKKOS_CLASS_LAMBDA(const int work_idx) {
        ^
/var/jenkins/workspace/Kokkos/core/src/setup/Kokkos_Setup_Cuda.hpp:87:34: note: expanded from macro 'KOKKOS_CLASS_LAMBDA'
#define KOKKOS_CLASS_LAMBDA [ =, *this ] __host__ __device__
                                 ^

Add RAJA check for accessing global variables in a RAJA loop

This would be a check to see if a RAJA lambda is trying to access a global variable. These variables are not captured by the lambda, and can cause problems when utilizing multiple memory spaces, since the lambda won't move the data over automatically.

Example:

static int my_static_var = 3;

int myfunction(double* data, int n)
{
int my_local_var = 5;
RAJA::for_all (RAJA::RangeSegment(0, n, [=](int i) {
...
data[i] += my_static_var; // catch this global variable access
data[i] += my_local_var; // this is fine. Lambda captured local variable
...
} );
}

Mock Kokkos::parallel functions

Generate mocked versions of Kokkos::parallel_for, Kokkos::parallel_reduce, and Kokkos::parallel_scan

Automate a process to rebase on upstream

I think we would like to rebase relatively often on upstream, but it would be a lot easier to do that if I automate that process.

Reproduction-Instructions For LLVM-CI

MinGW

• Download mingw-llvm for windows.
• install cmake for windows
• install git for windows (got it from visual studio)
• install python for windows
• set PATH=%PATH%;"c:\Program Files (x86)\Microsoft Visual Studio\Shared\Python37_64"
• set PATH=%PATH%;"f:\MinGW\llvm-mingw-20230614-msvcrt-x86_64\llvm-mingw-20230614-msvcrt-x86_64\bin"
• configure llvm:

cmake -S F:\MinGW\llvm-project\runtimes -B build-mingw -DLIBCXX_ENABLE_WERROR=YES -DLIBCXXABI_ENABLE_WERROR=YES -DLIBUNWIND_ENABLE_WERROR=YES -DLLVM_ENABLE_RUNTIMES="libcxx;libcxxabi;libunwind" -DLIBCXX_CXX_ABI=libcxxabi  -Clibcxx/cmake/caches/MinGW.cmake -DCMAKE_C_COMPILER=x86_64-w64-mingw32-clang -DCMAKE_CXX_COMPILER=x86_64-w64-mingw32-clang.exe -DCMAKE_ASM_COMPILER=x86_64-w64-mingw32-clang -G "MinGW Makefiles"

• mingw32-make
• python bin\llvm-lit.py -a ..\libcxx\test\std\containers\views\mdspan\mdspan

Port A bunch of checkers useful in some RAJA / CHAI/ CARE codes from clang-query to clang-tidy

let comment "### Disable default name for errors"
set bind-root false

let comment "### Identify device lambda contexts"
let b1 callExpr(hasArgument(0,hasType(asString("struct care::gpu"))))
let b2 callExpr(hasArgument(0,hasType(asString("struct care::parallel"))))
let b3 cxxMemberCallExpr(on(hasType(pointsTo(cxxRecordDecl(hasName("LoopFuser"))))))
let careHostDevicePtrCallee callee(cxxMethodDecl(ofClass(classTemplateSpecializationDecl(matchesName("care::host_device_ptr")))))
let rajaDeviceContext callExpr(anyOf(b1,b2,b3))
let inDeviceLambda hasAncestor(lambdaExpr(hasAncestor(rajaDeviceContext)))
let comment "### Identify code outside lambda contexts"
let notInsideLambda unless(hasAncestor(lambdaExpr()))
let comment "### Ignore implicit casts (to prevent duplicate matches)"
let notInImplicitCast unless(hasAncestor(implicitCastExpr()))
let hasCareHostDevPtrType hasType(qualType(hasCanonicalType(hasDeclaration(namedDecl(matchesName("care::host_device_ptr"))))))

######## USEFUL FOR RAJA USERS ###################

let comment "### Set up detection of assignment of pointer to 0 or NULL (instead of nullptr), except for with explicit casts"
match stmt(isExpansionInMainFile(), castExpr(hasCastKind("NullToPointer"), unless(hasAncestor(explicitCastExpr())), unless(hasSourceExpression(cxxNullPtrLiteralExpr())))).bind("assignment_of_pointer_to_0")

let comment "### Set up detection of implicit this within device RAJA context"
match cxxThisExpr(inDeviceLambda, notInImplicitCast).bind("implicit_use_of_this_inside_lambda_function")


#### USEFUL FOR CARE USERS (https://github.com/LLNL/CARE) ##################

let comment "### Set up detection of host_device_ptr operator[] outside RAJA context"
match cxxOperatorCallExpr(careHostDevicePtrCallee, notInsideLambda, hasOverloadedOperatorName("[]")).bind("use_of_host_device_ptr_operator[]_outside_lambda_function")

let comment "### Set up detection of host_device_ptr operator[] from pointer within inside lambda context (as a stand-in for RAJA contexts)"
let careHostDevPtrArrow stmt(allOf(memberExpr(isArrow()), expr(hasCareHostDevPtrType)))
match cxxOperatorCallExpr(careHostDevicePtrCallee, hasAncestor(lambdaExpr()), hasOverloadedOperatorName("[]"), hasDescendant(careHostDevPtrArrow)).bind("use_of_host_device_ptr_operator[]_from_pointer_inside_lambda_function")

let comment "### Set up detection of dereference of host_device_ptr, ignoring system headers (from vector<host_device_ptr>) and explicit casts"
match stmt(unaryOperator(unless(hasDescendant(explicitCastExpr())), unless(isExpansionInSystemHeader()), allOf(hasOperatorName("*"), hasUnaryOperand(hasDescendant(expr(hasCareHostDevPtrType)))))).bind("dereference_host_device_ptr")

let comment "### Set up detection of KeyValueSorter functions outside RAJA context"
let k1 cxxMethodDecl(hasName("key"))
let k2 cxxMethodDecl(hasName("setKey"))
let k3 cxxMethodDecl(hasName("value"))
let k4 cxxMethodDecl(hasName("setValue"))
let KeyValueSorterRAJACalls cxxMethodDecl(anyOf(k1,k2,k3,k4))
match cxxMemberCallExpr(on(hasType(cxxRecordDecl(hasName("KeyValueSorter")))), callee(KeyValueSorterRAJACalls), notInsideLambda).bind("use_of_KeyValueSort_RAJA_function_outside_lambda_function")

let comment "### Set up detection of assignment of raw pointer to host_device_ptr"
let comment "### This is searching for the constructor with bool as the third argument"
let comment "### Ignore when the raw pointer is explicitly cast"
match cxxConstructExpr(hasCareHostDevPtrType, hasArgument(2, hasType(asString("_Bool"))), unless(hasArgument(0, explicitCastExpr()))).bind("construction_of_host_device_ptr_from_raw_pointer")

############# USEFUL FOR CHAI USERS ########################

let comment "### Set up detection of reference of raw pointer in device lambda"
let comment "### ignoringImplicit and no implicitCastExpr ancestor is required to prevent duplicate matches"
match expr(inDeviceLambda, notInImplicitCast, ignoringImplicit(declRefExpr(to(varDecl(hasType(isAnyPointer()), unless(inDeviceLambda)))))).bind("capture_of_raw_pointer_in_lambda")

Add Kokkos specific analysis definition

This will allow us to put specific things into the Kokkos project that are protected by __kokkos_analyzer__. For example:

#if defined(__kokkos_analyzer__)
#define KOKKOS_FUNCTION ... __attribute((annotate("KOKKOS_FUNCTION")))
#endif

Currently we use __clang_analyzer__ for this, but that doesn't get turned on when clangd is running preventing certain checks from happening in the editor.

Ensure KOKKOS_FUNCTION on functions called in a parallel context

This check will ensure that functions called inside Kokkos::parallel_* functors and functions annotated with KOKKOS_FUNCTION are annotated with KOKKOS_FUNCTION.

At a minimum should detect the calls to oops in both not_oops and the KOKKOS_LAMBDA, this check will be made significantly easier if new versions of Kokkos are required due to an annotation attribute added to KOKKOS_FUNCTION.

int oops(int x){ return x;}

KOKKOS_FUNCTION
int not_oops(int x){return oops(x);}

Kokkos::parallel_for(7, KOKKOS_LAMBA(int x){printf("%d", oops(x));});

Also we need to be sure not to flag functions such as printf