Skip to content

Commit 60ec4e7

Browse files
zkoopmansgvisor-bot
zkoopmans
authored and
gvisor-bot
committed
Add 12.8 CUDA tests
Add 12.8 basic CUDA image and smoke tests. A later submission will add CUDA tests similar to the current 12.2 tests. PiperOrigin-RevId: 751593396
1 parent 27f3984 commit 60ec4e7

File tree

9 files changed

+534
-2
lines changed

9 files changed

+534
-2
lines changed

‎Makefile

+1-1
Original file line numberDiff line numberDiff line change
@@ -291,7 +291,7 @@ simple-tests: unit-tests # Compatibility target.
291291
.PHONY: simple-tests
292292

293293
# Images needed for GPU smoke tests.
294-
gpu-smoke-images: load-gpu_cuda-tests
294+
gpu-smoke-images: load-gpu_cuda-tests load-gpu_cuda-tests-12-8
295295
.PHONY: gpu-smoke-images
296296

297297
gpu-smoke-tests: gpu-smoke-images $(RUNTIME_BIN)

‎images/gpu/cuda-tests-12-8/Dockerfile

+49
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
FROM nvidia/cuda:12.8.1-devel-ubuntu22.04
2+
3+
WORKDIR /
4+
ENV PATH=$PATH:/usr/local/nvidia/bin
5+
RUN export DEBIAN_FRONTEND=noninteractive; \
6+
apt-get update && \
7+
apt-get install -y \
8+
build-essential \
9+
cmake \
10+
freeglut3 freeglut3-dev \
11+
git \
12+
golang \
13+
imagemagick \
14+
libegl-dev \
15+
libfreeimage3 libfreeimage-dev \
16+
libfreeimageplus3 libfreeimageplus-dev \
17+
libgles2-mesa-dev \
18+
libglfw3 libglfw3-dev \
19+
libglu1-mesa libglu1-mesa-dev \
20+
libxi-dev \
21+
libxmu-dev \
22+
llvm \
23+
mpich \
24+
pkg-config \
25+
vim \
26+
x11-xserver-utils \
27+
xdotool \
28+
xvfb \
29+
zlib1g zlib1g-dev
30+
31+
RUN git clone \
32+
https://github.com/NVIDIA/cuda-samples.git /cuda-samples && cd /cuda-samples && \
33+
git checkout 7b60178984e96bc09d066077d5455df71fee2a9f && cd /
34+
35+
RUN apt install -y wget && apt -y purge golang*
36+
37+
RUN wget https://go.dev/dl/go1.24.1.linux-amd64.tar.gz && tar -C /usr/local -xzf go1.24.1.linux-amd64.tar.gz && \
38+
ln -s /usr/local/go/bin/go /usr/local/bin/go
39+
40+
ADD *.cu *.h *.sh *.py *.cc /
41+
42+
RUN chmod 555 /*.sh && gcc -o /unsupported_ioctl /unsupported_ioctl.cc
43+
44+
RUN mkdir /cuda-samples/build && cd /cuda-samples/build && \
45+
cmake ..
46+
47+
# Override entrypoint to nothing, otherwise all invocations will have
48+
# a copyright notice printed, which breaks parsing the stdout logs.
49+
ENTRYPOINT []

‎images/gpu/cuda-tests-12-8/cuda_malloc.cu

+205
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,205 @@
1+
// Copyright 2023 The gVisor Authors.
2+
//
3+
// Licensed under the Apache License, Version 2.0 (the "License");
4+
// you may not use this file except in compliance with the License.
5+
// You may obtain a copy of the License at
6+
//
7+
// http://www.apache.org/licenses/LICENSE-2.0
8+
//
9+
// Unless required by applicable law or agreed to in writing, software
10+
// distributed under the License is distributed on an "AS IS" BASIS,
11+
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
// See the License for the specific language governing permissions and
13+
// limitations under the License.
14+
15+
#include <cuda_runtime.h>
16+
#include <err.h>
17+
#include <errno.h>
18+
#include <stdlib.h>
19+
#include <unistd.h>
20+
21+
#include <cstdint>
22+
#include <iostream>
23+
#include <random>
24+
25+
#include "cuda_test_util.h" // NOLINT(build/include)
26+
27+
__global__ void addKernel(std::uint32_t* data) {
28+
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
29+
data[index] += static_cast<std::uint32_t>(index);
30+
}
31+
32+
void TestMallocManagedRoundTrip(int device, unsigned int malloc_flags,
33+
bool prefetch) {
34+
constexpr size_t kNumBlocks = 32;
35+
constexpr size_t kNumThreads = 64;
36+
constexpr size_t kNumElems = kNumBlocks * kNumThreads;
37+
38+
std::uint32_t* data = nullptr;
39+
constexpr size_t kNumBytes = kNumElems * sizeof(*data);
40+
CHECK_CUDA(cudaMallocManaged(&data, kNumBytes, malloc_flags));
41+
42+
// Initialize all elements in the array with a random value on the host.
43+
std::default_random_engine rd;
44+
const std::uint32_t init_val =
45+
std::uniform_int_distribution<std::uint32_t>()(rd);
46+
for (size_t i = 0; i < kNumElems; i++) {
47+
data[i] = init_val;
48+
}
49+
50+
if (prefetch) {
51+
CHECK_CUDA(cudaMemPrefetchAsync(data, kNumBytes, device));
52+
}
53+
54+
// Mutate the array on the device.
55+
addKernel<<<kNumBlocks, kNumThreads>>>(data);
56+
CHECK_CUDA(cudaGetLastError());
57+
CHECK_CUDA(cudaDeviceSynchronize());
58+
59+
if (prefetch) {
60+
CHECK_CUDA(cudaMemPrefetchAsync(data, kNumBytes, cudaCpuDeviceId));
61+
}
62+
63+
// Check that the array has the expected result.
64+
for (size_t i = 0; i < kNumElems; i++) {
65+
std::uint32_t want = init_val + static_cast<std::uint32_t>(i);
66+
if (data[i] != want) {
67+
std::cout << "data[" << i << "]: got " << data[i] << ", wanted " << want
68+
<< " = " << init_val << " + " << i << std::endl;
69+
abort();
70+
}
71+
}
72+
73+
CHECK_CUDA(cudaFree(data));
74+
}
75+
76+
void TestMallocManagedReadWrite(int device) {
77+
constexpr size_t kNumBlocks = 32;
78+
constexpr size_t kNumThreads = 64;
79+
constexpr size_t kNumElems = kNumBlocks * kNumThreads;
80+
81+
std::uint32_t* data = nullptr;
82+
constexpr size_t kNumBytes = kNumElems * sizeof(*data);
83+
CHECK_CUDA(cudaMallocManaged(&data, kNumBytes, cudaMemAttachGlobal));
84+
85+
// Initialize all elements in the array with a random value on the host.
86+
std::default_random_engine rd;
87+
const std::uint32_t init_val =
88+
std::uniform_int_distribution<std::uint32_t>()(rd);
89+
for (size_t i = 0; i < kNumElems; i++) {
90+
data[i] = init_val;
91+
}
92+
93+
// Write the array's contents to a temporary file.
94+
char filename[] = "/tmp/cudaMallocManagedTest.XXXXXX";
95+
int fd = mkstemp(filename);
96+
if (fd < 0) {
97+
err(1, "mkstemp");
98+
}
99+
size_t done = 0;
100+
while (done < kNumBytes) {
101+
ssize_t n = write(fd, reinterpret_cast<char*>(data) + done,
102+
kNumBytes - done);
103+
if (n >= 0) {
104+
done += n;
105+
} else if (n < 0 && errno != EINTR) {
106+
err(1, "write");
107+
}
108+
}
109+
110+
// Mutate the array on the device.
111+
addKernel<<<kNumBlocks, kNumThreads>>>(data);
112+
CHECK_CUDA(cudaGetLastError());
113+
CHECK_CUDA(cudaDeviceSynchronize());
114+
115+
// Check that the array has the expected result.
116+
for (size_t i = 0; i < kNumElems; i++) {
117+
std::uint32_t want = init_val + static_cast<std::uint32_t>(i);
118+
if (data[i] != want) {
119+
std::cout << "data[" << i << "]: got " << data[i] << ", wanted " << want
120+
<< " = " << init_val << " + " << i << std::endl;
121+
abort();
122+
}
123+
}
124+
125+
// Read the array's original contents back from the temporary file.
126+
if (lseek(fd, 0, SEEK_SET) < 0) {
127+
err(1, "lseek");
128+
}
129+
done = 0;
130+
while (done < kNumBytes) {
131+
ssize_t n = read(fd, reinterpret_cast<char*>(data) + done,
132+
kNumBytes - done);
133+
if (n > 0) {
134+
done += n;
135+
} else if (n == 0) {
136+
errx(1, "read: unexpected EOF after %zu bytes", done);
137+
} else if (n < 0 && errno != EINTR) {
138+
err(1, "read");
139+
}
140+
}
141+
142+
// Check that the array matches what we originally wrote.
143+
for (size_t i = 0; i < kNumElems; i++) {
144+
std::uint32_t want = init_val;
145+
if (data[i] != want) {
146+
std::cout << "data[" << i << "]: got " << data[i] << ", wanted " << want
147+
<< " = " << init_val << " + " << i << std::endl;
148+
abort();
149+
}
150+
}
151+
152+
// Mutate the array on the device again.
153+
addKernel<<<kNumBlocks, kNumThreads>>>(data);
154+
CHECK_CUDA(cudaGetLastError());
155+
CHECK_CUDA(cudaDeviceSynchronize());
156+
157+
// Check that the array has the expected result again.
158+
for (size_t i = 0; i < kNumElems; i++) {
159+
std::uint32_t want = init_val + static_cast<std::uint32_t>(i);
160+
if (data[i] != want) {
161+
std::cout << "data[" << i << "]: got " << data[i] << ", wanted " << want
162+
<< " = " << init_val << " + " << i << std::endl;
163+
abort();
164+
}
165+
}
166+
167+
close(fd);
168+
CHECK_CUDA(cudaFree(data));
169+
}
170+
171+
int main() {
172+
int device;
173+
CHECK_CUDA(cudaGetDevice(&device));
174+
175+
std::cout << "Testing cudaMallocManaged(flags=cudaMemAttachGlobal)"
176+
<< std::endl;
177+
TestMallocManagedRoundTrip(device, cudaMemAttachGlobal, false);
178+
179+
int cma = 0;
180+
CHECK_CUDA(
181+
cudaDeviceGetAttribute(&cma, cudaDevAttrConcurrentManagedAccess, device));
182+
if (!cma) {
183+
std::cout << "cudaDevAttrConcurrentManagedAccess not available"
184+
<< std::endl;
185+
} else {
186+
std::cout << "Testing cudaMallocManaged(flags=cudaMemAttachGlobal) "
187+
"with prefetching"
188+
<< std::endl;
189+
TestMallocManagedRoundTrip(device, cudaMemAttachGlobal, true);
190+
std::cout << "Testing cudaMallocManaged(flags=cudaMemAttachHost)"
191+
<< std::endl;
192+
TestMallocManagedRoundTrip(device, cudaMemAttachHost, false);
193+
std::cout << "Testing cudaMallocManaged(flags=cudaMemAttachHost) "
194+
"with prefetching"
195+
<< std::endl;
196+
TestMallocManagedRoundTrip(device, cudaMemAttachHost, true);
197+
}
198+
199+
std::cout << "Testing read/write syscalls on cudaMallocManaged memory"
200+
<< std::endl;
201+
TestMallocManagedReadWrite(device);
202+
203+
std::cout << "All tests passed" << std::endl;
204+
return 0;
205+
}

‎images/gpu/cuda-tests-12-8/cuda_test_util.h

+42
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
// Copyright 2023 The gVisor Authors.
2+
//
3+
// Licensed under the Apache License, Version 2.0 (the "License");
4+
// you may not use this file except in compliance with the License.
5+
// You may obtain a copy of the License at
6+
//
7+
// http://www.apache.org/licenses/LICENSE-2.0
8+
//
9+
// Unless required by applicable law or agreed to in writing, software
10+
// distributed under the License is distributed on an "AS IS" BASIS,
11+
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
// See the License for the specific language governing permissions and
13+
// limitations under the License.
14+
15+
#ifndef THIRD_PARTY_GVISOR_IMAGES_GPU_CUDA_TESTS_CUDA_TEST_UTIL_H_
16+
#define THIRD_PARTY_GVISOR_IMAGES_GPU_CUDA_TESTS_CUDA_TEST_UTIL_H_
17+
18+
#include <iostream>
19+
20+
// cudaError_t is returned by CUDA runtime functions.
21+
#define CHECK_CUDA(expr) \
22+
do { \
23+
cudaError_t code = (expr); \
24+
if (code != cudaSuccess) { \
25+
std::cout << "Check failed at " << __FILE__ << ":" << __LINE__ << ": " \
26+
<< #expr << ": " << cudaGetErrorString(code) << std::endl; \
27+
abort(); \
28+
} \
29+
} while (0)
30+
31+
// CUresult is returned by CUDA driver functions.
32+
#define CHECK_CUDA_RESULT(expr) \
33+
do { \
34+
CUresult code = (expr); \
35+
if (code != CUDA_SUCCESS) { \
36+
std::cout << "Check failed at " << __FILE__ << ":" << __LINE__ << ": " \
37+
<< #expr << ": " << code << std::endl; \
38+
abort(); \
39+
} \
40+
} while (0)
41+
42+
#endif // THIRD_PARTY_GVISOR_IMAGES_GPU_CUDA_TESTS_CUDA_TEST_UTIL_H_

‎images/gpu/cuda-tests-12-8/run_cuda_test.py

+56
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
# Copyright 2025 The gVisor Authors.
2+
#
3+
# Licensed under the Apache License, Version 2.0 (the "License");
4+
# you may not use this file except in compliance with the License.
5+
# You may obtain a copy of the License at
6+
7+
# http://www.apache.org/licenses/LICENSE-2.0
8+
#
9+
# Unless required by applicable law or agreed to in writing, software
10+
# distributed under the License is distributed on an "AS IS" BASIS,
11+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
# See the License for the specific language governing permissions and
13+
# limitations under the License.
14+
15+
import argparse
16+
import os
17+
import shlex
18+
import subprocess
19+
import sys
20+
21+
ROOT_CMAKE_DIR = '/cuda-samples/build'
22+
SAMPLES_DIR = ROOT_CMAKE_DIR + '/Samples'
23+
24+
parser = argparse.ArgumentParser()
25+
26+
parser.add_argument(
27+
'test',
28+
help=(
29+
'Test to run. This should be some thing like'
30+
' "0_Introduction/UnifiedMemoryStreams"'
31+
),
32+
type=str,
33+
)
34+
35+
36+
def run_test(test_dir: str = ''):
37+
make_path = os.path.join(SAMPLES_DIR, test_dir)
38+
cmd = shlex.split(f'make -C {make_path}')
39+
subprocess.run(
40+
args=cmd,
41+
check=True,
42+
stdout=sys.stdout,
43+
stderr=sys.stderr,
44+
)
45+
test_path = [os.path.join(make_path, os.path.basename(test_dir))]
46+
subprocess.run(
47+
args=test_path,
48+
check=True,
49+
stdout=sys.stdout,
50+
stderr=sys.stderr,
51+
)
52+
53+
54+
if __name__ == '__main__':
55+
args = parser.parse_args(sys.argv[1:])
56+
run_test(test_dir=args.test)

‎images/gpu/cuda-tests-12-8/run_smoke.sh

+22
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
#!/bin/bash
2+
3+
# Copyright 2021 The gVisor Authors.
4+
#
5+
# Licensed under the Apache License, Version 2.0 (the "License");
6+
# you may not use this file except in compliance with the License.
7+
# You may obtain a copy of the License at
8+
#
9+
# http://www.apache.org/licenses/LICENSE-2.0
10+
#
11+
# Unless required by applicable law or agreed to in writing, software
12+
# distributed under the License is distributed on an "AS IS" BASIS,
13+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14+
# See the License for the specific language governing permissions and
15+
# limitations under the License.
16+
17+
set -euxo pipefail
18+
19+
cd /
20+
nvcc cuda_malloc.cu -o cuda_malloc -Wno-deprecated-gpu-targets
21+
./cuda_malloc
22+
echo 'All tests passed'

0 commit comments

Comments
 (0)