8000 Validate NVRTC binary output capacities before native writes by RicheyWorks · Pull Request #67 · jcuda/jcuda · GitHub
[go: up one dir, main page]
Skip to content

Validate NVRTC binary output capacities before native writes #67

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

Draft
RicheyWorks wants to merge 1 commit into
base: master
Choose a base branch
from
Draft

Validate NVRTC binary output capacities before native writes #67

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
8 changes: 8 additions & 0 deletions JCudaJava/src/main/java/jcuda/nvrtc/JNvrtc.java
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -318,6 +318,8 @@ public static int nvrtcGetCUBINSize(nvrtcProgram prog, long cubinSizeRet[])
* NVRTC_ERROR_INVALID_PROGRAM
*
* @see JNvrtc#nvrtcGetCUBINSize
* @throws IllegalArgumentException if the output array is smaller than
* the size reported by {@link #nvrtcGetCUBINSize}
*/
public static int nvrtcGetCUBIN(nvrtcProgram prog, byte cubin[])
{
Expand Down Expand Up @@ -357,6 +359,8 @@ public static int nvrtcGetNVVMSize(nvrtcProgram prog, long nvvmSizeRet[])
* NVRTC_SUCCESS
* NVRTC_ERROR_INVALID_INPUT
* NVRTC_ERROR_INVALID_PROGRAM
* @throws IllegalArgumentException if the output array is smaller than
* the size reported by {@link #nvrtcGetNVVMSize}
*/
public static int nvrtcGetNVVM(nvrtcProgram prog, byte nvvm[])
{
Expand Down Expand Up @@ -397,6 +401,8 @@ public static int nvrtcGetLTOIRSize(nvrtcProgram prog, long LTOIRSizeRet[])
* NVRTC_SUCCESS
* NVRTC_ERROR_INVALID_INPUT
* NVRTC_ERROR_INVALID_PROGRAM
* @throws IllegalArgumentException if the output array is smaller than
* the size reported by {@link #nvrtcGetLTOIRSize}
*/
public static int nvrtcGetLTOIR(nvrtcProgram prog, byte LTOIR[])
{
Expand Down Expand Up @@ -436,6 +442,8 @@ public static int nvrtcGetOptiXIRSize(nvrtcProgram prog, long optixirSizeRet[])
* NVRTC_SUCCESS
* NVRTC_ERROR_INVALID_INPUT
* NVRTC_ERROR_INVALID_PROGRAM
* @throws IllegalArgumentException if the output array is smaller than
* the size reported by {@link #nvrtcGetOptiXIRSize}
*/
public static int nvrtcGetOptiXIR(nvrtcProgram prog, byte optixir[])
{
Expand Down
140 changes: 140 additions & 0 deletions JCudaJava/src/test/java/jcuda/test/JNvrtcBinaryCapacityTest.java
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,140 @@
/* JCuda - Java bindings for CUDA - http://www.jcuda.org */
package jcuda.test;

import static jcuda.nvrtc.JNvrtc.*;
import static jcuda.nvrtc.nvrtcResult.*;
import static org.junit.Assert.*;
import java.io.File;
import java.nio.charset.StandardCharsets;
import java.nio.file.Files;
import java.util.ArrayList;
import java.util.Arrays;
import java.util.Collection;
import java.util.concurrent.TimeUnit;
import org.junit.Test;
import org.junit.runner.RunWith;
import org.junit.runners.Parameterized;
import jcuda.nvrtc.nvrtcProgram;

/** CPU-only output bounds checks; generated code is never loaded or launched. */
@RunWith(Parameterized.class)
public class JNvrtcBinaryCapacityTest
{
@Parameterized.Parameters(name = "{0}/{1}")
public static Collection<Object[]> cases()
{
Collection<Object[]> cases = new ArrayList<Object[]>();
for (String format : new String[] {"cubin", "nvvm", "lto", "optix"})
for (String mode : new String[] {"short", "empty", "exact", "padded", "invalid", "null"})
cases.add(new Object[] {format, mode});
return cases;
}
private final String format;
private final String mode;
public JNvrtcBinaryCapacityTest(String format, String mode)
{
this.format = format;
this.mode = mode;
}

@Test
public void respectsJavaArrayCapacity() throws Exception
{
String executable = System.getProperty("os.name").startsWith("Windows") ? "java.exe" : "java";
String java = new File(new File(System.getProperty("java.home"), "bin"), executable).getPath();
File log = File.createTempFile("jcuda-binary-capacity-", ".log");
Process process = new ProcessBuilder(java, "-Xcheck:jni",
"-Djava.library.path=" + System.getProperty("java.library.path"), "-cp",
System.getProperty("surefire.test.class.path", System.getProperty("java.class.path")),
getClass().getName(), format, mode).redirectErrorStream(true).redirectOutput(log).start();
try
{
assertTrue("NVRTC child timed out", process.waitFor(30, TimeUnit.SECONDS));
String output = new String(Files.readAllBytes(log.toPath()), StandardCharsets.UTF_8);
assertEquals(output, 0, process.exitValue());
assertTrue(output, output.contains("BINARY_CAPACITY_OK"));
}
finally
{
if (process.isAlive()) process.destroyForcibly();
log.delete();
}
}

private static int size(String format, nvrtcProgram p, long[] value)
{
if ("cubin".equals(format)) return nvrtcGetCUBINSize(p, value);
if ("nvvm".equals(format)) return nvrtcGetNVVMSize(p, value);
if ("lto".equals(format)) return nvrtcGetLTOIRSize(p, value);
return nvrtcGetOptiXIRSize(p, value);
}

private static int get(String format, nvrtcProgram p, byte[] value)
{
if ("cubin".equals(format)) return nvrtcGetCUBIN(p, value);
if ("nvvm".equals(format)) return nvrtcGetNVVM(p, value);
if ("lto".equals(format)) return nvrtcGetLTOIR(p, value);
return nvrtcGetOptiXIR(p, value);
}

public static void main(String[] args)
{
String format = args[0], mode = args[1];
nvrtcProgram program = new nvrtcProgram();
if ("invalid".equals(mode))
{
byte[] output = {42};
assertEquals(NVRTC_ERROR_INVALID_PROGRAM, get(format, program, output));
assertEquals(42, output[0]);
}
else if ("null".equals(mode))
{
assertThrows(NullPointerException.class, () -> get(format, program, null));
}
else
{
assertEquals(NVRTC_SUCCESS, nvrtcCreateProgram(program,
"extern \"C\" __global__ void foo(float *x) { x[0] += 1.0f; }",
"test.cu", 0, null, null));
try
{
String[] options = "cubin".equals(format)
? new String[] {"--gpu-architecture=sm_75"}
: new String[] {"--gpu-architecture=compute_75",
"optix".equals(format) ? "--optix-ir" : "--dlink-time-opt"};
int result = nvrtcCompileProgram(program, options.length, options);
if (result != NVRTC_SUCCESS)
{
String[] log = {null};
nvrtcGetProgramLog(program, log);
fail("Compilation: " + result + ": " + log[0]);
}
long[] required = {0};
assertEquals(NVRTC_SUCCESS, size(format, program, required));
assertTrue("Expected bounded positive compiler output", required[0] > 1 && required[0] < 1024 * 1024);
int length = (int)required[0];
if ("short".equals(mode) || "empty".equals(mode))
{
byte[] output = new byte["empty".equals(mode) ? 0 : length - 1];
Arrays.fill(output, (byte)42);
assertThrows(IllegalArgumentException.class, () -> get(format, program, output));
for (byte value : output) assertEquals(42, value);
}
else
{
byte[] expected = new byte[length];
assertEquals(NVRTC_SUCCESS, get(format, program, expected));
assertFalse("Compiler output must be nonzero", Arrays.equals(expected, new byte[length]));
byte[] output = new byte[length + ("padded".equals(mode) ? 17 : 0)];
Arrays.fill(output, (byte)42);
assertEquals(NVRTC_SUCCESS, get(format, program, output));
assertArrayEquals(expected, Arrays.copyOf(output, length));
for (int i = length; i < output.length; i++) assertEquals(42, output[i]);
}
System.out.println(format + " bytes=" + length);
}
finally { assertEquals(NVRTC_SUCCESS, nvrtcDestroyProgram(program)); }
}
System.out.println("BINARY_CAPACITY_OK");
}
}
48 changes: 44 additions & 4 deletions JNvrtcJNI/src/JNvrtc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -460,8 +460,18 @@ JNIEXPORT jint JNICALL Java_jcuda_nvrtc_JNvrtc_nvrtcGetCUBINNative

int length = env->GetArrayLength(cubin);
nvrtcProgram nativeProg = (nvrtcProgram)(intptr_t)getNativePointerValue(env, prog);
size_t requiredSize = 0;
int result = nvrtcGetCUBINSize(nativeProg, &requiredSize);
if (result != NVRTC_SUCCESS) return result;
if (requiredSize > (size_t)length)
{
ThrowByName(env, "java/lang/IllegalArgumentException",
"Output array is too small for nvrtcGetCUBIN");
return JCUDA_INTERNAL_ERROR;
}
char* nativeCubin = (char*)env->GetPrimitiveArrayCritical(cubin, NULL);
int result = nvrtcGetCUBIN(nativeProg, nativeCubin);
if (nativeCubin == NULL) return JCUDA_INTERNAL_ERROR;
result = nvrtcGetCUBIN(nativeProg, nativeCubin);
env->ReleasePrimitiveArrayCritical(cubin, nativeCubin, 0);
return result;

Expand Down Expand Up @@ -522,8 +532,18 @@ JNIEXPORT jint JNICALL Java_jcuda_nvrtc_JNvrtc_nvrtcGetNVVMNative

int length = env->GetArrayLength(nvvm);
nvrtcProgram nativeProg = (nvrtcProgram)(intptr_t)getNativePointerValue(env, prog);
size_t requiredSize = 0;
int result = nvrtcGetNVVMSize(nativeProg, &requiredSize);
if (result != NVRTC_SUCCESS) return result;
if (requiredSize > (size_t)length)
{
ThrowByName(env, "java/lang/IllegalArgumentException",
"Output array is too small for nvrtcGetNVVM");
return JCUDA_INTERNAL_ERROR;
}
char* nativeNvvm = (char*)env->GetPrimitiveArrayCritical(nvvm, NULL);
int result = nvrtcGetNVVM(nativeProg, nativeNvvm);
if (nativeNvvm == NULL) return JCUDA_INTERNAL_ERROR;
result = nvrtcGetNVVM(nativeProg, nativeNvvm);
env->ReleasePrimitiveArrayCritical(nvvm, nativeNvvm, 0);
return result;
}
Expand Down Expand Up @@ -585,8 +605,18 @@ JNIEXPORT jint JNICALL Java_jcuda_nvrtc_JNvrtc_nvrtcGetLTOIRNative

int length = env->GetArrayLength(LTOIR);
nvrtcProgram nativeProg = (nvrtcProgram)(intptr_t)getNativePointerValue(env, prog);
size_t requiredSize = 0;
int result = nvrtcGetLTOIRSize(nativeProg, &requiredSize);
if (result != NVRTC_SUCCESS) return result;
if (requiredSize > (size_t)length)
{
ThrowByName(env, "java/lang/IllegalArgumentException",
"Output array is too small for nvrtcGetLTOIR");
return JCUDA_INTERNAL_ERROR;
}
char* nativeLTOIR = (char*)env->GetPrimitiveArrayCritical(LTOIR, NULL);
int result = nvrtcGetLTOIR(nativeProg, nativeLTOIR);
if (nativeLTOIR == NULL) return JCUDA_INTERNAL_ERROR;
result = nvrtcGetLTOIR(nativeProg, nativeLTOIR);
env->ReleasePrimitiveArrayCritical(LTOIR, nativeLTOIR, 0);
return result;
}
Expand Down Expand Up @@ -647,8 +677,18 @@ JNIEXPORT jint JNICALL Java_jcuda_nvrtc_JNvrtc_nvrtcGetOptiXIRNative

int length = env->GetArrayLength(optixir);
nvrtcProgram nativeProg = (nvrtcProgram)(intptr_t)getNativePointerValue(env, prog);
size_t requiredSize = 0;
int result = nvrtcGetOptiXIRSize(nativeProg, &requiredSize);
if (result != NVRTC_SUCCESS) return result;
if (requiredSize > (size_t)length)
{
ThrowByName(env, "java/lang/IllegalArgumentException",
"Output array is too small for nvrtcGetOptiXIR");
return JCUDA_INTERNAL_ERROR;
}
char* nativeOptixir = (char*)env->GetPrimitiveArrayCritical(optixir, NULL);
int result = nvrtcGetOptiXIR(nativeProg, nativeOptixir);
if (nativeOptixir == NULL) return JCUDA_INTERNAL_ERROR;
result = nvrtcGetOptiXIR(nativeProg, nativeOptixir);
env->ReleasePrimitiveArrayCritical(optixir, nativeOptixir, 0);
return result;

Expand Down
0