Skip to content

Commit 269e98d

Browse files
jrpricejrprice
committed
Add support for LLVM 21
CI disabled for macOS due to issues with Homebrew version of LLVM. PCH disabled until LLVM 23 due to a bug with the relocatable PCH mechanism. Add support for llvm.experimental.noalias.scope.decl (just ignored). Fixes: #223
1 parent a1d42a9 commit 269e98d

File tree

10 files changed

+114
-6
lines changed

10 files changed

+114
-6
lines changed

‎.github/workflows/ci.yml‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@ on: [push, pull_request]
33

44
env:
55
RELEASE_VERSION: "21.10"
6-
RELEASE_LLVM: "20"
6+
RELEASE_LLVM: "21"
77

88
jobs:
99
check-format:
@@ -31,12 +31,14 @@ jobs:
3131
os: [ubuntu-22.04, macos-15]
3232
compiler: [gcc, clang]
3333
build_type: [Release, Debug]
34-
llvm: [18, 19, 20]
34+
llvm: [18, 19, 20, 21]
3535
exclude:
3636
- os: macos-15
3737
compiler: gcc
3838
- os: macos-15
3939
llvm: 20
40+
- os: macos-15
41+
llvm: 21
4042

4143
steps:
4244
- uses: actions/checkout@v2
@@ -105,7 +107,7 @@ jobs:
105107
fail-fast: false
106108
matrix:
107109
platform: [Win32, x64]
108-
llvm: [18, 19, 20]
110+
llvm: [18, 19, 20, 21]
109111

110112
steps:
111113
- uses: actions/checkout@v2

‎.github/workflows/install-deps.sh‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,12 @@ elif [[ "`uname`" == "MINGW64"* ]]; then
4747
curl -OL "$URL/$ARCHIVE"
4848
tar xf "$ARCHIVE" --strip-components 1 -C llvm-${LLVM_VERSION}/cmake
4949

50+
# Get third_party
51+
mkdir -p llvm-${LLVM_VERSION}/third-party
52+
ARCHIVE="third-party-${LLVM_FULL_VERSION}.src.tar.xz"
53+
curl -OL "$URL/$ARCHIVE"
54+
tar xf "$ARCHIVE" --strip-components 1 -C llvm-${LLVM_VERSION}/third-party
55+
5056
# Build LLVM + Clang
5157
mkdir -p llvm-${LLVM_VERSION}/build
5258
cd llvm-${LLVM_VERSION}/build

‎src/core/Program.cpp‎

Lines changed: 21 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -399,6 +399,7 @@ bool Program::build(BuildType buildType, const char* options,
399399
// Pre-compiled header
400400
char* pchdir = NULL;
401401
char* pch = NULL;
402+
#if LLVM_VERSION < 210 || LLVM_VERSION >= 230
402403
if (!checkEnv("OCLGRIND_DISABLE_PCH") &&
403404
(!strcmp(clstd, "-cl-std=CL1.2") || !strcmp(clstd, "-cl-std=CL2.0")))
404405
{
@@ -466,6 +467,7 @@ bool Program::build(BuildType buildType, const char* options,
466467
buildLog << "WARNING: Unable to determine precompiled header path\n";
467468
}
468469
}
470+
#endif
469471

470472
if (pch)
471473
{
@@ -490,6 +492,20 @@ bool Program::build(BuildType buildType, const char* options,
490492
clang::DiagnosticOptions* diagOpts = new clang::DiagnosticOptions();
491493
llvm::IntrusiveRefCntPtr<clang::DiagnosticIDs> diagID(
492494
new clang::DiagnosticIDs());
495+
496+
#if LLVM_VERSION >= 210
497+
clang::TextDiagnosticPrinter* diagConsumer =
498+
new clang::TextDiagnosticPrinter(buildLog, *diagOpts, false);
499+
clang::DiagnosticsEngine diags(diagID, *diagOpts, diagConsumer);
500+
501+
// Create compiler invocation and instance.
502+
std::shared_ptr<clang::CompilerInvocation> invocation(
503+
new clang::CompilerInvocation);
504+
clang::CompilerInvocation::CreateFromArgs(*invocation, args, diags);
505+
clang::CompilerInstance compiler(invocation);
506+
compiler.createDiagnostics(*llvm::vfs::getRealFileSystem(), diagConsumer,
507+
false);
508+
#else
493509
clang::TextDiagnosticPrinter* diagConsumer =
494510
new clang::TextDiagnosticPrinter(buildLog, diagOpts);
495511
clang::DiagnosticsEngine diags(diagID, diagOpts, diagConsumer);
@@ -509,6 +525,7 @@ bool Program::build(BuildType buildType, const char* options,
509525
clang::CompilerInvocation::CreateFromArgs(*invocation, args,
510526
compiler.getDiagnostics());
511527
compiler.setInvocation(invocation);
528+
#endif
512529

513530
// Remap include files
514531
std::unique_ptr<llvm::MemoryBuffer> buffer;
@@ -961,8 +978,10 @@ llvm::GetElementPtrInst* createScalarGEP(llvm::Value* ptr, llvm::Value* index,
961978
auto byteOffset =
962979
getTypeSize(store->getValueOperand()->getType()->getScalarType());
963980
auto* mul = llvm::BinaryOperator::CreateMul(
964-
index, llvm::Constant::getIntegerValue(index->getType(),
965-
llvm::APInt(32, byteOffset)));
981+
index,
982+
llvm::Constant::getIntegerValue(
983+
index->getType(),
984+
llvm::APInt(index->getType()->getIntegerBitWidth(), byteOffset)));
966985
mul->insertBefore(store);
967986
indices.push_back(mul);
968987
}

‎src/core/WorkItemBuiltins.cpp‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3557,6 +3557,11 @@ class WorkItemBuiltins
35573557
// TODO: Implement?
35583558
}
35593559

3560+
DEFINE_BUILTIN(llvm_noalias_scope_decl)
3561+
{
3562+
// TODO: Implement?
3563+
}
3564+
35603565
DEFINE_BUILTIN(llvm_memcpy)
35613566
{
35623567
const llvm::MemCpyInst* memcpyInst = (const llvm::MemCpyInst*)callInst;
@@ -3866,6 +3871,8 @@ BuiltinFunctionMap WorkItemBuiltins::initBuiltins()
38663871
ADD_PREFIX_BUILTIN("llvm.fabs.f", f1arg, F1ARG(fabs));
38673872
ADD_PREFIX_BUILTIN("llvm.lifetime.start", llvm_lifetime_start, NULL);
38683873
ADD_PREFIX_BUILTIN("llvm.lifetime.end", llvm_lifetime_end, NULL);
3874+
ADD_PREFIX_BUILTIN("llvm.experimental.noalias.scope.decl",
3875+
llvm_noalias_scope_decl, NULL);
38693876
ADD_PREFIX_BUILTIN("llvm.memcpy", llvm_memcpy, NULL);
38703877
ADD_PREFIX_BUILTIN("llvm.memmove", llvm_memcpy, NULL);
38713878
ADD_PREFIX_BUILTIN("llvm.memset", llvm_memset, NULL);

‎src/plugins/Uninitialized.cpp‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -999,6 +999,9 @@ void Uninitialized::handleIntrinsicInstruction(const WorkItem* workItem,
999999
case llvm::Intrinsic::lifetime_start:
10001000
// Do nothing
10011001
break;
1002+
case llvm::Intrinsic::experimental_noalias_scope_decl:
1003+
// Do nothing
1004+
break;
10021005
default:
10031006
FATAL_ERROR("Unsupported intrinsic %s",
10041007
llvm::Intrinsic::getName(I->getIntrinsicID()).data());
@@ -1640,7 +1643,11 @@ void Uninitialized::instructionExecuted(const WorkItem* workItem,
16401643
for (unsigned i = 0; i < newShadow.num; i++)
16411644
{
16421645
int index = shuffleInst->getMaskValue(i);
1646+
#if LLVM_VERSION >= 210
1647+
if (index == llvm::PoisonMaskElem)
1648+
#else
16431649
if (index == llvm::Value::UndefValueVal)
1650+
#endif
16441651
{
16451652
// Undef value are poisoned
16461653
memcpy(newShadow.data + i * newShadow.size, pv.data, newShadow.size);

‎tests/kernels/CMakeLists.txt‎

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ endforeach(${test})
2222
set_tests_properties(${KERNEL_TESTS} PROPERTIES
2323
ENVIRONMENT "OCLGRIND_PCH_DIR=${CMAKE_BINARY_DIR}/include/oclgrind")
2424

25-
# https://github.com/jrprice/Oclgrind/218
25+
# https://github.com/jrprice/Oclgrind/issues/218
2626
if ("${CMAKE_SYSTEM_NAME}" STREQUAL "Windows" AND
2727
${LLVM_VERSION_MAJOR} VERSION_EQUAL "18")
2828
set(XFAIL ${XFAIL} memcheck/static_array_padded_struct)
@@ -33,5 +33,11 @@ if (${LLVM_PACKAGE_VERSION} VERSION_LESS "19.0")
3333
set(XFAIL ${XFAIL} interactive/struct_member)
3434
endif()
3535

36+
# https://github.com/jrprice/Oclgrind/issues/221
37+
if (${LLVM_PACKAGE_VERSION} VERSION_GREATER_EQUAL "21.0")
38+
set(XFAIL ${XFAIL} interactive/pointers)
39+
set(XFAIL ${XFAIL} interactive/struct_member)
40+
endif()
41+
3642
# Expected failures
3743
set_tests_properties(${XFAIL} PROPERTIES WILL_FAIL TRUE)

‎tests/kernels/TESTS‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -33,6 +33,7 @@ bugs/llvm_memcpyopt_bug
3333
bugs/many_alloca
3434
bugs/multidim_array_in_struct
3535
bugs/null_argument
36+
bugs/restrict
3637
bugs/rhadd_overflow
3738
bugs/sroa_addrspace_cast
3839
bugs/write_vector_write_only_fp

‎tests/kernels/bugs/restrict.cl‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
kernel void reduce(uint n,
2+
global uint * restrict data,
3+
global uint * restrict result,
4+
local uint * restrict localData)
5+
{
6+
uint gid = get_global_id(0);
7+
uint lid = get_local_id(0);
8+
uint gsz = get_global_size(0);
9+
uint lsz = get_local_size(0);
10+
uint grp = get_group_id(0);
11+
12+
uint sum = 0;
13+
for (uint i = gid; i < n; i+=gsz)
14+
{
15+
sum += data[i];
16+
}
17+
18+
localData[lid] = sum;
19+
for (uint offset = lsz/2; offset > 0; offset/=2)
20+
{
21+
barrier(CLK_LOCAL_MEM_FENCE);
22+
if (lid < offset)
23+
{
24+
localData[lid] += localData[lid + offset];
25+
}
26+
}
27+
28+
if (lid == 0)
29+
{
30+
result[grp] = localData[lid];
31+
}
32+
}

‎tests/kernels/bugs/restrict.ref‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
EXACT Argument 'result': 64 bytes
2+
EXACT result[0] = 1560
3+
EXACT result[1] = 1624
4+
EXACT result[2] = 1688
5+
EXACT result[3] = 1752
6+
EXACT result[4] = 1816
7+
EXACT result[5] = 1880
8+
EXACT result[6] = 1944
9+
EXACT result[7] = 2008
10+
EXACT result[8] = 2072
11+
EXACT result[9] = 2136
12+
EXACT result[10] = 2200
13+
EXACT result[11] = 2264
14+
EXACT result[12] = 2328
15+
EXACT result[13] = 2392
16+
EXACT result[14] = 2456
17+
EXACT result[15] = 2520

‎tests/kernels/bugs/restrict.sim‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
restrict.cl
2+
reduce
3+
64 1 1
4+
4 1 1
5+
6+
<size=4>
7+
256
8+
9+
<size=1024 range=0:1:255>
10+
<size=64 fill=0 dump>
11+
<size=16>

0 commit comments

Comments
 (0)