Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Improve the capture of fatal cuda error #10884

Merged
merged 12 commits into from
Jun 7, 2022
Merged

Improve the capture of fatal cuda error #10884

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
823b145
update
sperlingxx May 18, 2022
404ce92
revert style
sperlingxx May 18, 2022
ef6e881
udpate
sperlingxx May 19, 2022
3dcb354
udpate
sperlingxx May 19, 2022
e0d026f
udpate
sperlingxx May 19, 2022
b610762
udpate
sperlingxx May 19, 2022
7e66695
update
sperlingxx May 20, 2022
b91e248
update
sperlingxx May 20, 2022
ac595a2
update
sperlingxx May 26, 2022
61058af
remove lines
sperlingxx May 26, 2022
3af3cc8
Merge remote-tracking branch 'origin/branch-22.06' into catch_cuda_fa…
sperlingxx May 26, 2022
182c15d
Merge remote-tracking branch 'origin/branch-22.08' into catch_cuda_fa…
sperlingxx May 30, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 1 addition & 2 deletions cpp/include/cudf/utilities/error.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,9 +114,8 @@ namespace detail {

inline void throw_cuda_error(cudaError_t error, const char* file, unsigned int line)
{
// Calls cudaGetLastError twice. It is nearly certain that a fatal error occurred if the second
// Calls cudaGetLastError again. It is nearly certain that a fatal error occurred if the second
// call doesn't return with cudaSuccess.
cudaGetLastError();
auto const last = cudaGetLastError();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Wait, what? The two calls are necessary to detect a fatal error vs a non-fatal. The first call clears any pending error state. If the second call still sees an error, then it's extremely likely that a sticky error has occurred.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

While the test case of fatal error can only be passed when I remove this line.

TEST(FatalCase, CudaFatalError)
{
  auto type = cudf::data_type{cudf::type_id::INT32};
  auto cv   = cudf::column_view(type, 256, (void*)256);
  cudf::binary_operation(cv, cv, cudf::binary_operator::ADD, type);
  EXPECT_THROW(CUDF_CUDA_TRY(cudaDeviceSynchronize()), cudf::fatal_cuda_error);
}

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agree that this seems like a dubious change. What specifically fails with the test without this change? Does it throw too early, in cudf::binary_operation or not throw at all? If the error is truly fatal, I don't see how removing a cudaGetLastError call is going to help this test pass. With a fatal error, we should be able to call cudaGetLastError as many times as we want, and it will never clear.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Then that means the error being generated isn't a true sticky error, which is admittedly surprising with the test you're doing here I'd expect an illegal access error (which I'm pretty sure is sticky).

You could condense this down a bit by just launching a kernel like:

__global__ void fatal_kernel() {
    __assert_fail(nullptr,nullptr,0,nullptr);
}
...
TEST(FatalCase, CudaFatalError)
{
  fatal_kernel<<<1,1>>>();
  EXPECT_THROW(CUDF_CUDA_TRY(cudaDeviceSynchronize()), cudf::fatal_cuda_error);
}

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Though you'd need to do this in a death test because it will corrupt the process context and leave the GPU unusable. See

cudf/cpp/tests/error/error_handling_test.cu

Lines 89 to 111 in 1db83e3

__global__ void assert_false_kernel() { cudf_assert(false && "this kernel should die"); }
__global__ void assert_true_kernel() { cudf_assert(true && "this kernel should live"); }
TEST(DebugAssertDeathTest, cudf_assert_false)
{
testing::FLAGS_gtest_death_test_style = "threadsafe";
auto call_kernel = []() {
assert_false_kernel<<<1, 1>>>();
// Kernel should fail with `cudaErrorAssert`
// This error invalidates the current device context, so we need to kill
// the current process. Running with EXPECT_DEATH spawns a new process for
// each attempted kernel launch
if (cudaErrorAssert == cudaDeviceSynchronize()) { std::abort(); }
// If we reach this point, the cudf_assert didn't work so we exit normally, which will cause
// EXPECT_DEATH to fail.
};
EXPECT_DEATH(call_kernel(), "this kernel should die");
}
for example.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I tried with fatal_kernel, but it throwed nothing.

I'm still not comfortable with removing a line of code without understanding why we're removing it. It may have helped your test case, but we need to understand how that was significant for that test (is it a problem with the test?) or how removing this will not create problems in other scenarios trying to detect fatal errors. If the CUDA error truly is fatal, it should not matter if we read the error an extra time. It should make it even more likely it truly is a fatal error if the error persists despite extra attempts at clearing it.

Copy link
Contributor Author

@sperlingxx sperlingxx May 23, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hi @jlowe, according to the description in CUDA doc: Returns the last error that has been produced by any of the runtime calls in the same host thread and resets it to cudaSuccess, the cudaGetLastError API works like popping the top error from the CUDA error stack ?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hi @jlowe, according to the description in CUDA doc: Returns the last error that has been produced by any of the runtime calls in the same host thread and resets it to cudaSuccess, the cudaGetLastError API works like popping the top error from the CUDA error stack ?

Yes, normally it clears the error, but there are categories of errors that are unclearable. These are the fatal errors we are trying to detect here. If you're finding that cudaGetLastError is able to clear an error then it seems that error is not actually a fatal error and we should not report it as such.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@sperlingxx see https://stackoverflow.com/questions/31642520/states-of-memory-data-after-cuda-exceptions/31642573#31642573

Copy link
Contributor Author

@sperlingxx sperlingxx May 24, 2022 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hi @jrhemstad, thank you for the link. However, my simple test suggests the second call of cudaGetLastError cleans up fatal errors as well, if I don't misunderstand anything.

  // a valid CUDA call
  int* p0;
  EXPECT_EQ(cudaMalloc(&p0, 128), cudaSuccess);

  // produce an unrecoverable CUDA error: cudaErrorIllegalAddress
  auto type = cudf::data_type{cudf::type_id::INT32};
  auto cv   = cudf::column_view(type, 256, (void*)256);
  cudf::binary_operation(cv, cv, cudf::binary_operator::ADD, type);
  // wait the illegal binary operation to finish, then capture the CUDA status
  EXPECT_EQ(cudaDeviceSynchronize(), cudaErrorIllegalAddress);
  EXPECT_EQ(cudaGetLastError(), cudaErrorIllegalAddress);
  EXPECT_EQ(cudaGetLastError(), cudaSuccess); // the second call returns success

  // Any subsequent CUDA calls will fail, since the CUDA context has been corrupted.
  int* p1;
  EXPECT_EQ(cudaMalloc(&p1, 128), cudaErrorIllegalAddress);
  EXPECT_EQ(cudaGetLastError(), cudaErrorIllegalAddress);
  EXPECT_EQ(cudaGetLastError(), cudaSuccess); // the second call returns success

  int* p2;
  EXPECT_EQ(cudaMalloc(&p2, 128), cudaErrorIllegalAddress);
  EXPECT_EQ(cudaGetLastError(), cudaErrorIllegalAddress);
  EXPECT_EQ(cudaGetLastError(), cudaSuccess); // the second call returns success

auto const msg = std::string{"CUDA error encountered at: " + std::string{file} + ":" +
std::to_string(line) + ": " + std::to_string(error) + " " +
Expand Down
19 changes: 16 additions & 3 deletions cpp/tests/error/error_handling_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,12 +16,13 @@

#include <cudf_test/base_fixture.hpp>

#include <cudf/binaryop.hpp>
#include <cudf/column/column_view.hpp>
#include <cudf/filling.hpp>
#include <cudf/utilities/error.hpp>

#include <rmm/cuda_stream.hpp>

#include <cstring>

TEST(ExpectsTest, FalseCondition)
{
EXPECT_THROW(CUDF_EXPECTS(false, "condition is false"), cudf::logic_error);
Expand Down Expand Up @@ -118,11 +119,23 @@ TEST(DebugAssert, cudf_assert_true)

#endif

TEST(FatalCase, CudaFatalError)
{
auto type = cudf::data_type{cudf::type_id::INT32};
auto cv = cudf::column_view(type, 256, (void*)256);
cudf::binary_operation(cv, cv, cudf::binary_operator::ADD, type);
EXPECT_THROW(CUDF_CUDA_TRY(cudaDeviceSynchronize()), cudf::fatal_cuda_error);
}

// These tests don't use CUDF_TEST_PROGRAM_MAIN because :
// 1.) They don't need the RMM Pool
// 2.) The RMM Pool interferes with the death test
// 3.) The order of test cases matters
int main(int argc, char** argv)
{
::testing::InitGoogleTest(&argc, argv);
return RUN_ALL_TESTS();
::testing::GTEST_FLAG(filter) = "-FatalCase.*";
int ret = RUN_ALL_TESTS();
::testing::GTEST_FLAG(filter) = "FatalCase.*";
return ret + RUN_ALL_TESTS();
}
123 changes: 76 additions & 47 deletions java/pom.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -136,7 +136,7 @@
<groupId>org.apache.arrow</groupId>
<artifactId>arrow-vector</artifactId>
<version>${arrow.version}</version>
<scope>test</scope>
<scope>test</scope>
</dependency>
<dependency>
<groupId>org.apache.parquet</groupId>
Expand Down Expand Up @@ -198,8 +198,30 @@
<configuration>
<excludes>
<exclude>**/CuFileTest.java</exclude>
<exclude>**/CudaTest.java</exclude>
</excludes>
</configuration>
jlowe marked this conversation as resolved.
Show resolved Hide resolved
<executions>
<execution>
<id>main-tests</id>
<goals>
<goal>test</goal>
</goals>
</execution>
<execution>
<id>fatal-cuda-test</id>
jlowe marked this conversation as resolved.
Show resolved Hide resolved
<goals>
<goal>test</goal>
</goals>
<configuration>
<includes>
<include>**/CudaTest.java</include>
</includes>
<reuseForks>false</reuseForks>
<test>*/CudaTest.java</test>
</configuration>
</execution>
</executions>
</plugin>
</plugins>
</build>
Expand Down Expand Up @@ -279,7 +301,7 @@
<nexusUrl>https://oss.sonatype.org/</nexusUrl>
<autoReleaseAfterClose>false</autoReleaseAfterClose>
</configuration>
</plugin>
</plugin>
</plugins>
</build>
</profile>
Expand All @@ -288,16 +310,16 @@
<build>
<resources>
<resource>
<!-- Include the properties file to provide the build information. -->
<directory>${project.build.directory}/extra-resources</directory>
<filtering>true</filtering>
<!-- Include the properties file to provide the build information. -->
<directory>${project.build.directory}/extra-resources</directory>
<filtering>true</filtering>
</resource>
<resource>
<directory>${basedir}/..</directory>
<targetPath>META-INF</targetPath>
<includes>
<include>LICENSE</include>
</includes>
<directory>${basedir}/..</directory>
<targetPath>META-INF</targetPath>
<includes>
<include>LICENSE</include>
</includes>
</resource>
</resources>
<pluginManagement>
Expand Down Expand Up @@ -338,6 +360,12 @@
<artifactId>junit-jupiter-engine</artifactId>
<version>5.4.2</version>
</dependency>
<dependency>
<!-- to get around bug https://github.com/junit-team/junit5/issues/1367 -->
<groupId>org.apache.maven.surefire</groupId>
<artifactId>surefire-logger-api</artifactId>
<version>2.21.0</version>
</dependency>
</dependencies>
</plugin>
<plugin>
Expand Down Expand Up @@ -384,9 +412,9 @@
executable="cmake">
<arg value="${basedir}/src/main/native"/>
<arg line="${cmake.ccache.opts}"/>
<arg value="-DCUDA_STATIC_RUNTIME=${CUDA_STATIC_RUNTIME}" />
<arg value="-DPER_THREAD_DEFAULT_STREAM=${PER_THREAD_DEFAULT_STREAM}" />
<arg value="-DUSE_GDS=${USE_GDS}" />
<arg value="-DCUDA_STATIC_RUNTIME=${CUDA_STATIC_RUNTIME}"/>
<arg value="-DPER_THREAD_DEFAULT_STREAM=${PER_THREAD_DEFAULT_STREAM}"/>
<arg value="-DUSE_GDS=${USE_GDS}"/>
<arg value="-DCMAKE_CXX_FLAGS=${cxx.flags}"/>
<arg value="-DCMAKE_EXPORT_COMPILE_COMMANDS=${CMAKE_EXPORT_COMPILE_COMMANDS}"/>
<arg value="-DCUDF_CPP_BUILD_DIR=${CUDF_CPP_BUILD_DIR}"/>
Expand All @@ -403,9 +431,10 @@
<arg value="${parallel.level}"/>
</exec>
<mkdir dir="${project.build.directory}/extra-resources"/>
<exec executable="bash" output="${project.build.directory}/extra-resources/cudf-java-version-info.properties">
<arg value="${project.basedir}/buildscripts/build-info"/>
<arg value="${project.version}"/>
<exec executable="bash"
output="${project.build.directory}/extra-resources/cudf-java-version-info.properties">
<arg value="${project.basedir}/buildscripts/build-info"/>
<arg value="${project.version}"/>
</exec>
</tasks>
</configuration>
Expand All @@ -427,31 +456,31 @@
</goals>
<configuration>
<source>
def sout = new StringBuffer(), serr = new StringBuffer()
//This only works on linux
def proc = 'ldd ${native.build.path}/libcudfjni.so'.execute()
proc.consumeProcessOutput(sout, serr)
proc.waitForOrKill(10000)
def libcudf = ~/libcudf.*\\.so\\s+=>\\s+(.*)libcudf.*\\.so\\s+.*/
def cudfm = libcudf.matcher(sout)
if (cudfm.find()) {
pom.properties['native.cudf.path'] = cudfm.group(1)
} else {
fail("Could not find cudf as a dependency of libcudfjni out> $sout err> $serr")
}
def sout = new StringBuffer(), serr = new StringBuffer()
//This only works on linux
def proc = 'ldd ${native.build.path}/libcudfjni.so'.execute()
proc.consumeProcessOutput(sout, serr)
proc.waitForOrKill(10000)
def libcudf = ~/libcudf.*\\.so\\s+=>\\s+(.*)libcudf.*\\.so\\s+.*/
def cudfm = libcudf.matcher(sout)
if (cudfm.find()) {
pom.properties['native.cudf.path'] = cudfm.group(1)
} else {
fail("Could not find cudf as a dependency of libcudfjni out> $sout err> $serr")
}

def nvccout = new StringBuffer(), nvccerr = new StringBuffer()
def nvccproc = 'nvcc --version'.execute()
nvccproc.consumeProcessOutput(nvccout, nvccerr)
nvccproc.waitForOrKill(10000)
def cudaPattern = ~/Cuda compilation tools, release ([0-9]+)/
def cm = cudaPattern.matcher(nvccout)
if (cm.find()) {
def classifier = 'cuda' + cm.group(1)
pom.properties['cuda.classifier'] = classifier
} else {
fail('could not find CUDA version')
}
def nvccout = new StringBuffer(), nvccerr = new StringBuffer()
def nvccproc = 'nvcc --version'.execute()
nvccproc.consumeProcessOutput(nvccout, nvccerr)
nvccproc.waitForOrKill(10000)
def cudaPattern = ~/Cuda compilation tools, release ([0-9]+)/
def cm = cudaPattern.matcher(nvccout)
if (cm.find()) {
def classifier = 'cuda' + cm.group(1)
pom.properties['cuda.classifier'] = classifier
} else {
fail('could not find CUDA version')
}
</source>
</configuration>
</execution>
Expand Down Expand Up @@ -479,13 +508,13 @@
<groupId>org.apache.maven.plugins</groupId>
<artifactId>maven-surefire-plugin</artifactId>
<configuration>
<!-- you can turn this off, by passing -DtrimStackTrace=true when running tests -->
<trimStackTrace>false</trimStackTrace>
<redirectTestOutputToFile>true</redirectTestOutputToFile>
<systemPropertyVariables>
<ai.rapids.refcount.debug>${ai.rapids.refcount.debug}</ai.rapids.refcount.debug>
<ai.rapids.cudf.nvtx.enabled>${ai.rapids.cudf.nvtx.enabled}</ai.rapids.cudf.nvtx.enabled>
</systemPropertyVariables>
<!-- you can turn this off, by passing -DtrimStackTrace=true when running tests -->
<trimStackTrace>false</trimStackTrace>
<redirectTestOutputToFile>true</redirectTestOutputToFile>
<systemPropertyVariables>
<ai.rapids.refcount.debug>${ai.rapids.refcount.debug}</ai.rapids.refcount.debug>
<ai.rapids.cudf.nvtx.enabled>${ai.rapids.cudf.nvtx.enabled}</ai.rapids.cudf.nvtx.enabled>
</systemPropertyVariables>
</configuration>
</plugin>
<plugin>
Expand Down
23 changes: 17 additions & 6 deletions java/src/main/native/include/jni_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -807,14 +807,12 @@ inline void jni_cuda_check(JNIEnv *const env, cudaError_t cuda_status) {
}

// Throw a new exception only if one is not pending then always return with the specified value
#define JNI_CHECK_CUDA_ERROR(env, class_name, e, ret_val) \
#define JNI_CHECK_CUDA_ERROR(env, class_name, msg, e_code, ret_val) \
{ \
if (env->ExceptionOccurred()) { \
return ret_val; \
} \
std::string n_msg = e.what() == nullptr ? "" : e.what(); \
jstring j_msg = env->NewStringUTF(n_msg.c_str()); \
jint e_code = static_cast<jint>(e.error_code()); \
jlowe marked this conversation as resolved.
Show resolved Hide resolved
jstring j_msg = env->NewStringUTF(msg); \
jclass ex_class = env->FindClass(class_name); \
if (ex_class != NULL) { \
jmethodID ctor_id = env->GetMethodID(ex_class, "<init>", "(Ljava/lang/String;I)V"); \
Expand Down Expand Up @@ -856,12 +854,25 @@ inline void jni_cuda_check(JNIEnv *const env, cudaError_t cuda_status) {
JNI_CHECK_THROW_NEW(env, cudf::jni::OOM_CLASS, what.c_str(), ret_val); \
} \
catch (const cudf::fatal_cuda_error &e) { \
JNI_CHECK_CUDA_ERROR(env, cudf::jni::CUDA_FATAL_ERROR_CLASS, e, ret_val); \
const char *what = e.what() == nullptr ? "" : e.what(); \
auto e_code = static_cast<jint>(e.error_code()); \
JNI_CHECK_CUDA_ERROR(env, cudf::jni::CUDA_FATAL_ERROR_CLASS, what, e_code, ret_val); \
} \
catch (const cudf::cuda_error &e) { \
JNI_CHECK_CUDA_ERROR(env, cudf::jni::CUDA_ERROR_CLASS, e, ret_val); \
const char *what = e.what() == nullptr ? "" : e.what(); \
auto e_code = static_cast<jint>(e.error_code()); \
JNI_CHECK_CUDA_ERROR(env, cudf::jni::CUDA_ERROR_CLASS, what, e_code, ret_val); \
} \
catch (const std::exception &e) { \
/* Double check whether the thrown exception is unrecoverable CUDA error or not. */ \
/* Like cudf::detail::throw_cuda_error, it is nearly certain that a fatal error */ \
/* occurred if the second call doesn't return with cudaSuccess. */ \
auto const last = cudaDeviceSynchronize(); \
if (cudaSuccess != last && last == cudaGetLastError()) { \
jlowe marked this conversation as resolved.
Show resolved Hide resolved
const char *what = e.what() == nullptr ? "" : e.what(); \
auto code = static_cast<jint>(last); \
JNI_CHECK_CUDA_ERROR(env, cudf::jni::CUDA_FATAL_ERROR_CLASS, what, code, ret_val); \
} \
/* If jni_exception caught then a Java exception is pending and this will not overwrite it. */ \
JNI_CHECK_THROW_NEW(env, class_name, e.what(), ret_val); \
}
Expand Down
53 changes: 52 additions & 1 deletion java/src/test/java/ai/rapids/cudf/CudaTest.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,13 +16,15 @@

package ai.rapids.cudf;

import org.junit.jupiter.api.Test;
import org.junit.jupiter.api.*;

import static org.junit.jupiter.api.Assertions.*;

@TestMethodOrder(MethodOrderer.OrderAnnotation.class)
public class CudaTest {

@Test
@Order(1)
public void testGetCudaRuntimeInfo() {
// The driver version is not necessarily larger than runtime version. Drivers of previous
// version are also able to support runtime of later version, only if they support same
Expand All @@ -33,6 +35,7 @@ public void testGetCudaRuntimeInfo() {
}

@Test
@Order(2)
public void testCudaException() {
assertThrows(CudaException.class, () -> {
try {
Expand All @@ -44,5 +47,53 @@ public void testCudaException() {
}
}
);
// non-fatal CUDA error will not fail subsequent CUDA calls
try (ColumnVector cv = ColumnVector.fromBoxedInts(1, 2, 3, 4, 5)) {
}
}

@Test
@Order(3)
public void testCudaFatalException() {
jlowe marked this conversation as resolved.
Show resolved Hide resolved
try (ColumnView cv = ColumnView.fromDeviceBuffer(new BadDeviceBuffer(), 0, DType.INT8, 256);
ColumnView ret = cv.sub(cv);
HostColumnVector hcv = ret.copyToHost()) {
} catch (CudaException ignored) {
}

// CUDA API invoked by libcudf failed because of previous unrecoverable fatal error
assertThrows(CudaFatalException.class, () -> {
try (ColumnView cv = ColumnView.fromDeviceBuffer(new BadDeviceBuffer(), 0, DType.INT8, 256);
jlowe marked this conversation as resolved.
Show resolved Hide resolved
HostColumnVector hcv = cv.copyToHost()) {
} catch (CudaFatalException ex) {
assertEquals(CudaException.CudaError.cudaErrorIllegalAddress, ex.cudaError);
throw ex;
}
});
}

@Test
@Order(4)
public void testCudaFatalExceptionFromRMM() {
// CUDA API invoked by RMM failed because of previous unrecoverable fatal error
jlowe marked this conversation as resolved.
Show resolved Hide resolved
assertThrows(CudaFatalException.class, () -> {
try (ColumnVector cv = ColumnVector.fromBoxedInts(1, 2, 3, 4, 5)) {
} catch (CudaFatalException ex) {
assertEquals(CudaException.CudaError.cudaErrorIllegalAddress, ex.cudaError);
throw ex;
}
});
}

private static class BadDeviceBuffer extends BaseDeviceMemoryBuffer {
public BadDeviceBuffer() {
super(256L, 256L, (MemoryBufferCleaner) null);
}

@Override
public MemoryBuffer slice(long offset, long len) {
return null;
}
}

}