Author: Alex Bradbury Date: 2026-11-02T06:39:51Z New Revision: f7c2c4938886373b9fae1c4ea221f0db01835b74
URL: https://github.com/llvm/llvm-project/commit/f7c2c4938886373b9fae1c4ea221f0db01835b74 DIFF: https://github.com/llvm/llvm-project/commit/f7c2c4938886373b9fae1c4ea221f0db01835b74.diff LOG: [clang-repl] Remove host assumptions from Interpreter/CUDA tests (#237794) PR #236206 recently aadded tests in clang/test/Interpreter/CUDA that run on any Linux host with the NVPTX target and JIT support, but these tests have some assumptions that aren't true on some of the RISC-V buildbots: * The stub libcudart.so is linked with the default linker, which fails on hosts without one. This patch addresses this by linking with -fuse-ld=lld and adding `REQUIRES: lld`, as clang/test/Driver does. * The CHECK lines will break on targets such as RISC-V that add signext to i32 returns and arguments. I've adjusted the CHECK lines to handle this in a similar way to similar lines in Interpreter/emit-llvm.cpp. Debugging the buildbot issues and searching the repo for similar patterns in existing tests was LLM assisted. Added: Modified: clang/test/Interpreter/CUDA/device-function-codegen.cu clang/test/Interpreter/CUDA/device-function-template-codegen.cu clang/test/Interpreter/CUDA/device-module-verifier.cu clang/test/Interpreter/CUDA/emit-llvm.cu clang/test/Interpreter/CUDA/empty-device-module.cu clang/test/Interpreter/CUDA/host-and-device-codegen.cu clang/test/Interpreter/CUDA/lit.local.cfg clang/test/Interpreter/CUDA/memory-codegen.cu clang/test/Interpreter/CUDA/ptx-version.cu clang/test/Interpreter/CUDA/sanity-codegen.cu Removed: ################################################################################ diff --git a/clang/test/Interpreter/CUDA/device-function-codegen.cu b/clang/test/Interpreter/CUDA/device-function-codegen.cu index c6f9d72a09abc3..8b496b3c3c69a3 100644 --- a/clang/test/Interpreter/CUDA/device-function-codegen.cu +++ b/clang/test/Interpreter/CUDA/device-function-codegen.cu @@ -1,7 +1,8 @@ +// REQUIRES: lld // Tests __device__ function calls // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ @@ -18,7 +19,7 @@ __global__ void test_kernel(int* value) { test_device(value); } int* devptr = nullptr; cudaMalloc((void **) &devptr, sizeof(int)); test_kernel<<<1,1>>>(devptr); -// CHECK: call i32 @cudaConfigureCall( +// CHECK: call {{.*}}i32 @cudaConfigureCall( // CHECK: call void @_Z26__device_stub__test_kernelPi( %quit diff --git a/clang/test/Interpreter/CUDA/device-function-template-codegen.cu b/clang/test/Interpreter/CUDA/device-function-template-codegen.cu index 5e03e7cc909043..861fb765720f8b 100644 --- a/clang/test/Interpreter/CUDA/device-function-template-codegen.cu +++ b/clang/test/Interpreter/CUDA/device-function-template-codegen.cu @@ -1,7 +1,8 @@ +// REQUIRES: lld // Tests device function templates // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ @@ -17,7 +18,7 @@ __global__ void test_kernel(int* value) { *value = sum(40, 2); } int* devptr = nullptr; cudaMalloc((void **) &devptr, sizeof(int)); test_kernel<<<1,1>>>(devptr); -// CHECK: call i32 @cudaConfigureCall( +// CHECK: call {{.*}}i32 @cudaConfigureCall( // CHECK: call void @_Z26__device_stub__test_kernelPi( %quit diff --git a/clang/test/Interpreter/CUDA/device-module-verifier.cu b/clang/test/Interpreter/CUDA/device-module-verifier.cu index 33a33228a5a646..f18ac10de47e38 100644 --- a/clang/test/Interpreter/CUDA/device-module-verifier.cu +++ b/clang/test/Interpreter/CUDA/device-module-verifier.cu @@ -1,8 +1,9 @@ +// REQUIRES: lld // Release builds skip the IR verifier. Turn it on and check that the first // device module passes it. // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ diff --git a/clang/test/Interpreter/CUDA/emit-llvm.cu b/clang/test/Interpreter/CUDA/emit-llvm.cu index fa13055f67f18a..b687acb5c12053 100644 --- a/clang/test/Interpreter/CUDA/emit-llvm.cu +++ b/clang/test/Interpreter/CUDA/emit-llvm.cu @@ -1,7 +1,8 @@ +// REQUIRES: lld // -emit-llvm prints the device IR of each input before its host IR. // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ diff --git a/clang/test/Interpreter/CUDA/empty-device-module.cu b/clang/test/Interpreter/CUDA/empty-device-module.cu index a34c351a7174dc..b89e4b320c737a 100644 --- a/clang/test/Interpreter/CUDA/empty-device-module.cu +++ b/clang/test/Interpreter/CUDA/empty-device-module.cu @@ -1,9 +1,10 @@ +// REQUIRES: lld // Tests host-only inputs. They produce an empty device module, and emitting // PTX for it must not be reported as a failure just because no pass changed // the module. // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ diff --git a/clang/test/Interpreter/CUDA/host-and-device-codegen.cu b/clang/test/Interpreter/CUDA/host-and-device-codegen.cu index 0006750c124288..8f9356b0fcf829 100644 --- a/clang/test/Interpreter/CUDA/host-and-device-codegen.cu +++ b/clang/test/Interpreter/CUDA/host-and-device-codegen.cu @@ -1,7 +1,8 @@ +// REQUIRES: lld // Checks that a function is available in both __host__ and __device__ // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ @@ -15,13 +16,13 @@ __global__ void kernel(int * output){ *output = sum(40,2); } // CHECK: define {{.*}}void @_Z21__device_stub__kernelPi( int host_sum = sum(41, 1); -// CHECK: call {{.*}}i32 @_Z3sumii(i32 noundef 41, i32 noundef 1) +// CHECK: call {{.*}}i32 @_Z3sumii( // CHECK: define linkonce_odr {{.*}}i32 @_Z3sumii( int * deviceVar; cudaMalloc((void **) &deviceVar, sizeof(int)); kernel<<<1,1>>>(deviceVar); -// CHECK: call i32 @cudaConfigureCall( +// CHECK: call {{.*}}i32 @cudaConfigureCall( // CHECK: call void @_Z21__device_stub__kernelPi( %quit diff --git a/clang/test/Interpreter/CUDA/lit.local.cfg b/clang/test/Interpreter/CUDA/lit.local.cfg index f8ea448a2d6dc6..a918a1a7cab3ef 100644 --- a/clang/test/Interpreter/CUDA/lit.local.cfg +++ b/clang/test/Interpreter/CUDA/lit.local.cfg @@ -1,6 +1,7 @@ import subprocess import lit.util +from lit.llvm import llvm_config def have_host_clang_repl_cuda(): @@ -47,3 +48,6 @@ if ( elif have_host_clang_repl_cuda(): # The tests that execute a kernel carry REQUIRES: host-supports-cuda. config.available_features.add("host-supports-cuda") + +if llvm_config.use_lld(required=False): + config.available_features.add("lld") diff --git a/clang/test/Interpreter/CUDA/memory-codegen.cu b/clang/test/Interpreter/CUDA/memory-codegen.cu index 032e9e0b4f4f82..a405a84ff63f6d 100644 --- a/clang/test/Interpreter/CUDA/memory-codegen.cu +++ b/clang/test/Interpreter/CUDA/memory-codegen.cu @@ -1,7 +1,8 @@ +// REQUIRES: lld // Tests cudaMemcpy and writes from kernel // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ @@ -20,7 +21,7 @@ cudaMalloc((void **) &devptr, sizeof(int)); // CHECK: call {{.*}}@cudaMalloc( test_func<<<1,1>>>(devptr); -// CHECK: call i32 @cudaConfigureCall( +// CHECK: call {{.*}}i32 @cudaConfigureCall( // CHECK: call void @_Z24__device_stub__test_funcPi( cudaMemcpy(&var, devptr, sizeof(int), cudaMemcpyDeviceToHost); diff --git a/clang/test/Interpreter/CUDA/ptx-version.cu b/clang/test/Interpreter/CUDA/ptx-version.cu index 536fd37de458ea..d93a546ca08f55 100644 --- a/clang/test/Interpreter/CUDA/ptx-version.cu +++ b/clang/test/Interpreter/CUDA/ptx-version.cu @@ -1,7 +1,8 @@ +// REQUIRES: lld // The PTX version of the embedded fatbin must follow the pinned feature. // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ diff --git a/clang/test/Interpreter/CUDA/sanity-codegen.cu b/clang/test/Interpreter/CUDA/sanity-codegen.cu index f7c8bae1bdb3cf..0a70b4df0e33ed 100644 --- a/clang/test/Interpreter/CUDA/sanity-codegen.cu +++ b/clang/test/Interpreter/CUDA/sanity-codegen.cu @@ -1,6 +1,7 @@ +// REQUIRES: lld // RUN: rm -rf %t // RUN: mkdir -p %t/lib -// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared \ +// RUN: %clang --target=%host-jit-triple -nostdlib -fPIC -shared -fuse-ld=lld \ // RUN: %S/Inputs/cudart.c -o %t/lib/libcudart.so // RUN: cat %s | clang-repl --cuda --cuda-path=%t -Xcc -nocudainc \ // RUN: -Xcc -nocudalib -Xcc -include -Xcc %S/Inputs/cuda.h \ @@ -14,7 +15,7 @@ __global__ void test_func() {} // CHECK: call ptr @__cudaRegisterFatBinary( test_func<<<1,1>>>(); -// CHECK: call i32 @cudaConfigureCall( +// CHECK: call {{.*}}i32 @cudaConfigureCall( // CHECK: call void @_Z24__device_stub__test_funcv() %quit _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
