diff --git a/Notes.txt b/Notes.txt
new file mode 100644
index 0000000000000000000000000000000000000000..fd9aca631a72bdfaf2bfbcfb13a8662bf66d8d18
--- /dev/null
+++ b/Notes.txt
@@ -0,0 +1,3 @@
+Maybe we can use Java 7's method handles for this?
+http://java.sun.com/developer/technicalArticles/DynTypeLang/
+
diff --git a/build.xml b/build.xml
index 555f879c461d9a7559800a3585ea2c69a9004a76..77d91cf98a4f539c7f5a1eb3b25df59818051c64 100644
--- a/build.xml
+++ b/build.xml
@@ -1,6 +1,6 @@
<?xml version="1.0"?>
-<project name="aparapi" default="help" basedir=".">
+<project name="aparapi" default="build" basedir=".">
<property environment="env" />
<condition property="x86_or_x86_64" value="x86" else="x86_64"> <or><os arch="x86" /><os arch="i386"/></or> </condition>
diff --git a/com.amd.aparapi.jni/Aparapi.sln b/com.amd.aparapi.jni/Aparapi.sln
new file mode 100644
index 0000000000000000000000000000000000000000..3dc133919c83052b2af503ae0d977da4692e418f
--- /dev/null
+++ b/com.amd.aparapi.jni/Aparapi.sln
@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 11.00
+# Visual C++ Express 2010
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "Aparapi", "Aparapi.vcxproj", "{F0194884-D4F7-25AC-C03E-4D9229BB60CE}"
+EndProject
+Global
+ GlobalSection(SolutionConfigurationPlatforms) = preSolution
+ Debug|Win32 = Debug|Win32
+ Release|Win32 = Release|Win32
+ EndGlobalSection
+ GlobalSection(ProjectConfigurationPlatforms) = postSolution
+ {F0194884-D4F7-25AC-C03E-4D9229BB60CE}.Debug|Win32.ActiveCfg = Debug|Win32
+ {F0194884-D4F7-25AC-C03E-4D9229BB60CE}.Debug|Win32.Build.0 = Debug|Win32
+ {F0194884-D4F7-25AC-C03E-4D9229BB60CE}.Release|Win32.ActiveCfg = Release|Win32
+ {F0194884-D4F7-25AC-C03E-4D9229BB60CE}.Release|Win32.Build.0 = Release|Win32
+ EndGlobalSection
+ GlobalSection(SolutionProperties) = preSolution
+ HideSolutionNode = FALSE
+ EndGlobalSection
+EndGlobal
diff --git a/com.amd.aparapi.jni/Aparapi.suo b/com.amd.aparapi.jni/Aparapi.suo
new file mode 100644
index 0000000000000000000000000000000000000000..681b017555520c85b5df6002227c378c8cc38f37
Binary files /dev/null and b/com.amd.aparapi.jni/Aparapi.suo differ
diff --git a/com.amd.aparapi.jni/Aparapi.vcxproj b/com.amd.aparapi.jni/Aparapi.vcxproj
new file mode 100644
index 0000000000000000000000000000000000000000..61dfd198108364217b194c87c048421d7c0e6b4f
--- /dev/null
+++ b/com.amd.aparapi.jni/Aparapi.vcxproj
@@ -0,0 +1,90 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+ <ItemGroup Label="ProjectConfigurations">
+ <ProjectConfiguration Include="Debug|Win32">
+ <Configuration>Debug</Configuration>
+ <Platform>Win32</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release|Win32">
+ <Configuration>Release</Configuration>
+ <Platform>Win32</Platform>
+ </ProjectConfiguration>
+ </ItemGroup>
+ <PropertyGroup Label="Globals">
+ <Keyword>Win32Proj</Keyword>
+ </PropertyGroup>
+ <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
+ <ConfigurationType>DynamicLibrary</ConfigurationType>
+ <UseDebugLibraries>true</UseDebugLibraries>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
+ <ConfigurationType>DynamicLibrary</ConfigurationType>
+ <UseDebugLibraries>false</UseDebugLibraries>
+ </PropertyGroup>
+ <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+ <ImportGroup Label="ExtensionSettings">
+ </ImportGroup>
+ <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+ <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+ </ImportGroup>
+ <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+ <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+ </ImportGroup>
+ <PropertyGroup Label="UserMacros" />
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+ <LinkIncremental>true</LinkIncremental>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+ <LinkIncremental>true</LinkIncremental>
+ </PropertyGroup>
+ <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+ <ClCompile>
+ <PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;WIN32;_DEBUG;_WINDOWS;_USRDLL;APARAPI_EXPORTS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+ <AdditionalIncludeDirectories>C:\Users\Gary\aparapi.googlecode.com\branches\AddExtensionMechanism\com.amd.aparapi.jni\include;C:\Program Files\AMD APP\include;C:\Program Files\Java\jdk1.6.0_26\include;C:\Program Files\Java\jdk1.6.0_26\include\win32;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+ <RuntimeLibrary>MultiThreadedDebugDLL</RuntimeLibrary>
+ <WarningLevel>Level3</WarningLevel>
+ <DebugInformationFormat>ProgramDatabase</DebugInformationFormat>
+ <Optimization>Disabled</Optimization>
+ </ClCompile>
+ <Link>
+ <TargetMachine>MachineX86</TargetMachine>
+ <GenerateDebugInformation>true</GenerateDebugInformation>
+ <SubSystem>Windows</SubSystem>
+ <AdditionalLibraryDirectories>C:\Program Files\AMD APP\lib\x86;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
+ <AdditionalDependencies>kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;OpenCL.lib;%(AdditionalDependencies)</AdditionalDependencies>
+ <OutputFile>dist\aparapi_x86.dll</OutputFile>
+ </Link>
+ </ItemDefinitionGroup>
+ <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+ <ClCompile>
+ <PreprocessorDefinitions>WIN32;NDEBUG;_WINDOWS;_USRDLL;APARAPI_EXPORTS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+ <AdditionalIncludeDirectories>C:\Program Files\AMD APP\include\CL C:\Program Files\Java\jdk1.6.0_26\include C:\Program Files\Java\jdk1.6.0_26\include\x86;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+ <RuntimeLibrary>MultiThreadedDLL</RuntimeLibrary>
+ <WarningLevel>Level3</WarningLevel>
+ <DebugInformationFormat>ProgramDatabase</DebugInformationFormat>
+ </ClCompile>
+ <Link>
+ <TargetMachine>MachineX86</TargetMachine>
+ <GenerateDebugInformation>true</GenerateDebugInformation>
+ <SubSystem>Windows</SubSystem>
+ <EnableCOMDATFolding>true</EnableCOMDATFolding>
+ <OptimizeReferences>true</OptimizeReferences>
+ </Link>
+ </ItemDefinitionGroup>
+ <ItemGroup>
+ <ClCompile Include="src\cpp\aparapi.cpp" />
+ <ClCompile Include="src\cpp\jni.cpp" />
+ <ClCompile Include="src\cpp\jniHelper.cpp" />
+ </ItemGroup>
+ <ItemGroup>
+ <ClInclude Include="include\com_amd_aparapi_JNI.h" />
+ <ClInclude Include="include\com_amd_aparapi_KernelRunner.h" />
+ <ClInclude Include="src\cpp\aparapi.h" />
+ <ClInclude Include="src\cpp\common.h" />
+ <ClInclude Include="src\cpp\jniHelper.h" />
+ </ItemGroup>
+ <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+ <ImportGroup Label="ExtensionTargets">
+ </ImportGroup>
+</Project>
\ No newline at end of file
diff --git a/com.amd.aparapi.jni/Aparapi.vcxproj.filters b/com.amd.aparapi.jni/Aparapi.vcxproj.filters
new file mode 100644
index 0000000000000000000000000000000000000000..7e7d2b21a10c1eb5f54a352b363f03528876e3d6
--- /dev/null
+++ b/com.amd.aparapi.jni/Aparapi.vcxproj.filters
@@ -0,0 +1,45 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+ <ItemGroup>
+ <Filter Include="Source Files">
+ <UniqueIdentifier>{4FC737F1-C7A5-4376-A066-2A32D752A2FF}</UniqueIdentifier>
+ <Extensions>cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx</Extensions>
+ </Filter>
+ <Filter Include="Header Files">
+ <UniqueIdentifier>{93995380-89BD-4b04-88EB-625FBE52EBFB}</UniqueIdentifier>
+ <Extensions>h;hpp;hxx;hm;inl;inc;xsd</Extensions>
+ </Filter>
+ <Filter Include="Resource Files">
+ <UniqueIdentifier>{67DA6AB6-F800-4c08-8B7A-83BB121AAD01}</UniqueIdentifier>
+ <Extensions>rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav</Extensions>
+ </Filter>
+ </ItemGroup>
+ <ItemGroup>
+ <ClCompile Include="src\cpp\aparapi.cpp">
+ <Filter>Source Files</Filter>
+ </ClCompile>
+ <ClCompile Include="src\cpp\jni.cpp">
+ <Filter>Source Files</Filter>
+ </ClCompile>
+ <ClCompile Include="src\cpp\jniHelper.cpp">
+ <Filter>Source Files</Filter>
+ </ClCompile>
+ </ItemGroup>
+ <ItemGroup>
+ <ClInclude Include="include\com_amd_aparapi_JNI.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
+ <ClInclude Include="include\com_amd_aparapi_KernelRunner.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
+ <ClInclude Include="src\cpp\aparapi.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
+ <ClInclude Include="src\cpp\jniHelper.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
+ <ClInclude Include="src\cpp\common.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
+ </ItemGroup>
+</Project>
\ No newline at end of file
diff --git a/com.amd.aparapi.jni/Aparapi.vcxproj.user b/com.amd.aparapi.jni/Aparapi.vcxproj.user
new file mode 100644
index 0000000000000000000000000000000000000000..695b5c78b91edfc29f77823eb642fc9ead8e15f1
--- /dev/null
+++ b/com.amd.aparapi.jni/Aparapi.vcxproj.user
@@ -0,0 +1,3 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+</Project>
\ No newline at end of file
diff --git a/com.amd.aparapi.jni/build.xml b/com.amd.aparapi.jni/build.xml
index 7f0477a75969cb9eb14e0a55017f98e7c0184d19..334af523ef709ba10b5d7861eafa9e0f67cbef30 100644
--- a/com.amd.aparapi.jni/build.xml
+++ b/com.amd.aparapi.jni/build.xml
@@ -1,22 +1,22 @@
<?xml version="1.0"?>
<!--
-You should not have to edit this file
-First consider editing the properties in build.properties
+ You should not have to edit this file
+ First consider editing the properties in build.properties
-->
<project name="com.amd.aparapi.jni" default="build" basedir=".">
- <!--
- DO NOT EDIT BELOW THIS LINE
- -->
- <echo>OS Name: ${os.name}</echo>
- <echo>OS Version: ${os.version}</echo>
- <echo>OS Arch: ${os.arch}</echo>
-
- <property name="build.compiler" value="javac1.6"/>
- <property name="ant.build.javac.source" value="1.6"/>
- <property name="ant.build.javac.target" value="1.6"/>
-
+ <!--
+ DO NOT EDIT BELOW THIS LINE
+ -->
+ <echo>OS Name: ${os.name}</echo>
+ <echo>OS Version: ${os.version}</echo>
+ <echo>OS Arch: ${os.arch}</echo>
+
+ <property name="build.compiler" value="javac1.6"/>
+ <property name="ant.build.javac.source" value="1.6"/>
+ <property name="ant.build.javac.target" value="1.6"/>
+
<property environment="env" />
<!-- we need env.PATH for msvc only -->
@@ -264,9 +264,9 @@ First consider editing the properties in build.properties
<target name="clean" depends="check">
<delete dir="include" />
- <delete dir="${basedir}/dist" />
- <!-- Legacy cleanup -->
- <delete file="libaparapi_${x86_or_x86_64}.dylib" />
+ <delete dir="${basedir}/dist" />
+ <!-- Legacy cleanup -->
+ <delete file="libaparapi_${x86_or_x86_64}.dylib" />
<delete file="libaparapi_${x86_or_x86_64}.so" />
<delete file="aparapi_${x86_or_x86_64}.dll" />
<delete file="aparapi.dll" />
@@ -274,17 +274,24 @@ First consider editing the properties in build.properties
<delete file="aparapi.o" />
<delete file="aparapi.lib" />
<delete file="aparapi.exp" />
+ <delete file="opencljni.obj" />
+ <delete file="opencljni.o" />
+ <delete file="jniHelper.obj" />
+ <delete file="jniHelper.o" />
+ <delete file="clHelper.obj" />
+ <delete file="clHelper.o" />
</target>
<target name="javah">
<mkdir dir="include" />
<javah classpath="..\com.amd.aparapi\dist\aparapi.jar" destdir="include" force="true">
<class name="com.amd.aparapi.KernelRunner" />
+ <class name="com.amd.aparapi.OpenCLJNI" />
</javah>
</target>
<target name="gcc" if="use.gcc">
- <mkdir dir="${basedir}/dist"/>
+ <mkdir dir="${basedir}/dist"/>
<echo message="linuxcc ${os.arch}" />
<exec executable="g++">
<arg value="-O3" />
@@ -298,20 +305,20 @@ First consider editing the properties in build.properties
<arg value="-o" />
<arg value="${basedir}/dist/libaparapi_${x86_or_x86_64}.so" />
<arg value="src/cpp/aparapi.cpp" />
+ <arg value="src/cpp/opencljni.cpp" />
+ <arg value="src/cpp/jniHelper.cpp" />
+ <arg value="src/cpp/clHelper.cpp" />
<arg value="-L${amd.app.sdk.dir}/lib/${x86_or_x86_64}" />
<arg value="-lOpenCL" />
</exec>
</target>
<target name="gcc_mac" if="use.gcc_mac">
- <mkdir dir="${basedir}/dist"/>
- <echo message="clang++ ${os.arch}" />
- <exec executable="clang++">
+ <mkdir dir="${basedir}/dist"/>
+ <echo message="gcc ${os.arch}" />
+ <exec executable="g++">
<arg value="-O3" />
<arg value="-g" />
- <arg value="-m64" />
- <arg value="-Wno-format" />
- <arg value="-Wno-deprecated-writable-strings" />
<arg value="-fPIC" />
<arg value="-I/System/Library/Frameworks/JavaVM.framework/Headers" />
<arg value="-I/System/Library/Frameworks/OpenCL.framework/Headers" />
@@ -320,14 +327,17 @@ First consider editing the properties in build.properties
<arg value="-o" />
<arg value="${basedir}/dist/libaparapi_${x86_or_x86_64}.dylib" />
<arg value="src/cpp/aparapi.cpp" />
+ <arg value="src/cpp/opencljni.cpp" />
+ <arg value="src/cpp/jniHelper.cpp" />
+ <arg value="src/cpp/clHelper.cpp" />
<arg value="-framework" />
<arg value="OpenCL" />
</exec>
</target>
<target name="msvc" if="use.msvc">
- <mkdir dir="${basedir}\dist"/>
- <echo message="msvc ${os.arch}" />
+ <mkdir dir="${basedir}\dist"/>
+ <echo message="msvc ${os.arch}" />
<exec executable="${msvc.dir}\vc\bin\${optional.amd64.subdir}cl.exe">
<env key="PATH" path="${env.PATH};${msvc.dir}\\Common7\\IDE" />
<arg value="/nologo" />
@@ -340,6 +350,9 @@ First consider editing the properties in build.properties
<arg value="/Iinclude" />
<arg value="/I${amd.app.sdk.dir}\include" />
<arg value="src\cpp\aparapi.cpp" />
+ <arg value="src\cpp\opencljni.cpp" />
+ <arg value="src\cpp\jniHelper.cpp" />
+ <arg value="src/cpp/clHelper.cpp" />
<arg value="/LD" />
<arg value="/link" />
<arg value="/libpath:${msvc.dir}\vc\lib\${optional.amd64.subdir}" />
@@ -351,8 +364,8 @@ First consider editing the properties in build.properties
</target>
<target name="mingw" if="use.mingw">
- <mkdir dir="${basedir}\dist"/>
- <echo message="mingw ${os.arch}" />
+ <mkdir dir="${basedir}\dist"/>
+ <echo message="mingw ${os.arch}" />
<exec executable="${mingw.dir}/bin/g++">
<env key="PATH" path="${env.PATH};${mingw.dir}/bin" />
<arg value="-Wall" />
@@ -366,6 +379,9 @@ First consider editing the properties in build.properties
<arg value="-o" />
<arg value="${basedir}\dist\aparapi_${x86_or_x86_64}.dll" />
<arg value="src\cpp\aparapi.cpp" />
+ <arg value="src\cpp\opencljni.cpp" />
+ <arg value="src/cpp/jniHelper.cpp" />
+ <arg value="src/cpp/clHelper.cpp" />
<arg value="-L${amd.app.sdk.dir}\lib\${x86_or_x86_64}" />
<arg value="-lOpenCL" />
</exec>
diff --git a/com.amd.aparapi.jni/src/cpp/aparapi.cpp b/com.amd.aparapi.jni/src/cpp/aparapi.cpp
index 1a25846e35456915a8893e69734332e9de0f957e..6e89085a88af2c45d408167bd4df718e089c91ac 100644
--- a/com.amd.aparapi.jni/src/cpp/aparapi.cpp
+++ b/com.amd.aparapi.jni/src/cpp/aparapi.cpp
@@ -33,174 +33,16 @@
direct product is subject to national security controls as identified on the Commerce Control List (currently
found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
- and SecurityÂ’s website at http://www.bis.doc.gov/.
+ and Security?s website at http://www.bis.doc.gov/.
*/
-
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <time.h>
-
-#ifndef __APPLE__
-#include <malloc.h>
-#endif
-
-#include <sys/types.h>
-#ifndef _WIN32
-#include <unistd.h>
-#endif
-
-#ifndef __APPLE__
-#include <CL/cl.h>
-#else
-#include <cl.h>
-#endif
-
-#include <jni.h>
-
-#define JNIExceptionChecker(){\
- fprintf(stderr, "line %d\n", __LINE__);\
- if ((jenv)->ExceptionOccurred()) {\
- (jenv)->ExceptionDescribe(); /* write to console */\
- (jenv)->ExceptionClear();\
- }\
-}
-
-
-#if defined (_WIN32)
-#include "windows.h"
-#define alignedMalloc(size, alignment)\
- _aligned_malloc(size, alignment)
-#else
-#define alignedMalloc(size, alignment)\
- memalign(alignment, size)
-#endif
-
-
-class MicrosecondTimer{
-
-#if defined (_WIN32)
- private:
- __int64 freq;
- __int64 startValue;
- public:
- void start(){
- QueryPerformanceFrequency((LARGE_INTEGER*)&freq);
- QueryPerformanceCounter((LARGE_INTEGER*)&startValue);
- }
- void end(char *msg){
- __int64 endValue;
- QueryPerformanceCounter((LARGE_INTEGER*)&endValue);
- int us = (int)((endValue-startValue)* 1000000.0 / freq);
- fprintf(stderr, "%s=%d\n", msg, us);
- }
-
-#else
- public:
- void start(){
- }
- void end(char *msg){
- }
-
-#endif
-};
-
-MicrosecondTimer timer;
-
-
+#include "common.h"
+#include "jniHelper.h"
+#include "clHelper.h"
+#define APARAPI_SOURCE
+#include "aparapi.h"
#include "com_amd_aparapi_KernelRunner.h"
+#define JNI_JAVA(type, className, methodName) JNIEXPORT type JNICALL Java_com_amd_aparapi_##className##_##methodName
-#define CHECK_NO_RETURN(condition, msg) if(condition){\
- fprintf(stderr, "!!!!!!! %s failed !!!!!!!\n", msg);\
-}
-
-#define CHECK(condition, msg) if(condition){\
- fprintf(stderr, "!!!!!!! %s failed !!!!!!!\n", msg);\
- return 0;\
-}
-
-
-#define ASSERT_CL_NO_RETURN(msg) if (status != CL_SUCCESS){\
- fprintf(stderr, "!!!!!!! %s failed: %s\n", msg, CLErrString(status));\
-}
-
-#define ASSERT_CL(msg) if (status != CL_SUCCESS){\
- ASSERT_CL_NO_RETURN(msg)\
- return 0;\
-}
-
-#define PRINT_CL_ERR(status, msg) fprintf(stderr, "!!!!!!! %s failed %s\n", msg, CLErrString(status));
-
-#define ASSERT_FIELD(id) CHECK_NO_RETURN(id##FieldID == 0, "No such field as " #id)
-
-
-
-static const char *CLErrString(cl_int status) {
- static struct { cl_int code; const char *msg; } error_table[] = {
- { CL_SUCCESS, "success" },
- { CL_DEVICE_NOT_FOUND, "device not found", },
- { CL_DEVICE_NOT_AVAILABLE, "device not available", },
- { CL_COMPILER_NOT_AVAILABLE, "compiler not available", },
- { CL_MEM_OBJECT_ALLOCATION_FAILURE, "mem object allocation failure", },
- { CL_OUT_OF_RESOURCES, "out of resources", },
- { CL_OUT_OF_HOST_MEMORY, "out of host memory", },
- { CL_PROFILING_INFO_NOT_AVAILABLE, "profiling not available", },
- { CL_MEM_COPY_OVERLAP, "memcopy overlaps", },
- { CL_IMAGE_FORMAT_MISMATCH, "image format mismatch", },
- { CL_IMAGE_FORMAT_NOT_SUPPORTED, "image format not supported", },
- { CL_BUILD_PROGRAM_FAILURE, "build program failed", },
- { CL_MAP_FAILURE, "map failed", },
- { CL_INVALID_VALUE, "invalid value", },
- { CL_INVALID_DEVICE_TYPE, "invalid device type", },
- { CL_INVALID_PLATFORM, "invlaid platform",},
- { CL_INVALID_DEVICE, "invalid device",},
- { CL_INVALID_CONTEXT, "invalid context",},
- { CL_INVALID_QUEUE_PROPERTIES, "invalid queue properties",},
- { CL_INVALID_COMMAND_QUEUE, "invalid command queue",},
- { CL_INVALID_HOST_PTR, "invalid host ptr",},
- { CL_INVALID_MEM_OBJECT, "invalid mem object",},
- { CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, "invalid image format descriptor ",},
- { CL_INVALID_IMAGE_SIZE, "invalid image size",},
- { CL_INVALID_SAMPLER, "invalid sampler",},
- { CL_INVALID_BINARY, "invalid binary",},
- { CL_INVALID_BUILD_OPTIONS, "invalid build options",},
- { CL_INVALID_PROGRAM, "invalid program ",},
- { CL_INVALID_PROGRAM_EXECUTABLE, "invalid program executable",},
- { CL_INVALID_KERNEL_NAME, "invalid kernel name",},
- { CL_INVALID_KERNEL_DEFINITION, "invalid definition",},
- { CL_INVALID_KERNEL, "invalid kernel",},
- { CL_INVALID_ARG_INDEX, "invalid arg index",},
- { CL_INVALID_ARG_VALUE, "invalid arg value",},
- { CL_INVALID_ARG_SIZE, "invalid arg size",},
- { CL_INVALID_KERNEL_ARGS, "invalid kernel args",},
- { CL_INVALID_WORK_DIMENSION , "invalid work dimension",},
- { CL_INVALID_WORK_GROUP_SIZE, "invalid work group size",},
- { CL_INVALID_WORK_ITEM_SIZE, "invalid work item size",},
- { CL_INVALID_GLOBAL_OFFSET, "invalid global offset",},
- { CL_INVALID_EVENT_WAIT_LIST, "invalid event wait list",},
- { CL_INVALID_EVENT, "invalid event",},
- { CL_INVALID_OPERATION, "invalid operation",},
- { CL_INVALID_GL_OBJECT, "invalid gl object",},
- { CL_INVALID_BUFFER_SIZE, "invalid buffer size",},
- { CL_INVALID_MIP_LEVEL, "invalid mip level",},
- { CL_INVALID_GLOBAL_WORK_SIZE, "invalid global work size",},
- { 0, NULL },
- };
- static char unknown[25];
- int ii;
-
- for (ii = 0; error_table[ii].msg != NULL; ii++) {
- if (error_table[ii].code == status) {
- return error_table[ii].msg;
- }
- }
-#ifdef _WIN32
- _snprintf(unknown, sizeof unknown, "unknown error %d", status);
-#else
- snprintf(unknown, sizeof(unknown), "unknown error %d", status);
-#endif
- return unknown;
-}
class Range{
public:
static jclass rangeClazz;
@@ -356,7 +198,11 @@ class KernelArg{
isStatic = jenv->GetBooleanField(argObj, isStaticFieldID);
jstring nameString = (jstring)jenv->GetObjectField(argObj, nameFieldID);
const char *nameChars = jenv->GetStringUTFChars(nameString, NULL);
+#ifdef _WIN32
+ name=_strdup(nameChars);
+#else
name=strdup(nameChars);
+#endif
jenv->ReleaseStringUTFChars(nameString, nameChars);
}
@@ -696,7 +542,7 @@ void dispose(JNIEnv *jenv){
context = (cl_context)0;
}
if (commandQueues){
- for (int dev=0; dev<deviceIdc; dev++){
+ for (unsigned dev=0; dev<deviceIdc; dev++){
status = clReleaseCommandQueue((cl_command_queue)commandQueues[dev]);
ASSERT_CL_NO_RETURN("clReleaseCommandQueue()");
commandQueues[dev] = (cl_command_queue)0;
@@ -775,8 +621,8 @@ void unpinAll(JNIEnv* jenv) {
-
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_disposeJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
+JNI_JAVA(jint, KernelRunner, disposeJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext != NULL){
@@ -789,14 +635,14 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_disposeJNI(JNIEnv *jenv
void idump(char *str, void *ptr, int size){
int * iptr = (int *)ptr;
- for (int i=0; i<size/sizeof(int); i++){
+ for (unsigned i=0; i<size/sizeof(int); i++){
fprintf(stderr, "%s%4d %d\n", str, i, iptr[i]);
}
}
void fdump(char *str, void *ptr, int size){
float * fptr = (float *)ptr;
- for (int i=0; i<size/sizeof(float); i++){
+ for (unsigned i=0; i<size/sizeof(float); i++){
fprintf(stderr, "%s%4d %6.2f\n", str, i, fptr[i]);
}
}
@@ -968,8 +814,8 @@ jint updateNonPrimitiveReferences(JNIEnv *jenv, jobject jobj, JNIContext* jniCon
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *jenv,
- jobject jobj, jlong jniContextHandle, jobject _range, jboolean needSync, jint passes) {
+JNI_JAVA(jint, KernelRunner, runKernelJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jobject _range, jboolean needSync, jint passes) {
Range range(jenv, _range);
@@ -1004,9 +850,6 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *je
}
}
- if (jniContext->isVerbose()){
- timer.start();
- }
// Need to capture array refs
if (jniContext->firstRun || needSync) {
@@ -1259,7 +1102,7 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *je
jniContext->exec = new ProfileInfo[passes];
for (int passid=0; passid<passes; passid++){
- for (int dev =0; dev < jniContext->deviceIdc; dev++){ // this will always be 1 until we reserect multi-dim support
+ for (unsigned dev =0; dev < jniContext->deviceIdc; dev++){ // this will always be 1 until we reserect multi-dim support
//size_t offset = 1; // (size_t)((range.globalDims[0]/jniContext->deviceIdc)*dev);
status = clSetKernelArg(jniContext->kernel, kernelArgPos, sizeof(passid), &(passid));
if (status != CL_SUCCESS) {
@@ -1340,7 +1183,7 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *je
return status;
}
- for (int dev = 0; dev < jniContext->deviceIdc; dev++){
+ for (unsigned dev = 0; dev < jniContext->deviceIdc; dev++){
status = clReleaseEvent(jniContext->executeEvents[dev]);
if (status != CL_SUCCESS) {
PRINT_CL_ERR(status, "clReleaseEvent() read event");
@@ -1457,7 +1300,7 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *je
return executeStatus;
}
- for (int dev=0; dev<jniContext->deviceIdc; dev++){
+ for (unsigned int dev=0; dev<jniContext->deviceIdc; dev++){
status = clReleaseEvent(jniContext->executeEvents[dev]);
if (status != CL_SUCCESS) {
@@ -1486,9 +1329,6 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *je
}
jniContext->firstRun = false;
- if (jniContext->isVerbose()){
- timer.end("elapsed");
- }
//fprintf(stderr, "About to return %d from exec\n", status);
return(status);
@@ -1496,8 +1336,8 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_runKernelJNI(JNIEnv *je
// we return the JNIContext from here
-JNIEXPORT jlong JNICALL Java_com_amd_aparapi_KernelRunner_initJNI(JNIEnv *jenv, jclass clazz, jobject kernelObject,
- jint flags) {
+JNI_JAVA(jlong, KernelRunner, initJNI)
+ (JNIEnv *jenv, jclass clazz, jobject kernelObject, jint flags) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = new JNIContext(jenv, kernelObject, flags);
@@ -1510,41 +1350,18 @@ JNIEXPORT jlong JNICALL Java_com_amd_aparapi_KernelRunner_initJNI(JNIEnv *jenv,
}
-JNIEXPORT jlong JNICALL Java_com_amd_aparapi_KernelRunner_buildProgramJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jstring source) {
+JNI_JAVA(jlong, KernelRunner, buildProgramJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jstring source) {
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext == NULL){
return 0;
}
cl_int status = CL_SUCCESS;
- const char *sourceChars = jenv->GetStringUTFChars(source, NULL);
- CHECK(sourceChars == NULL, "jenv->GetStringUTFChars() returned null" );
-
- size_t sourceSize[] = { strlen(sourceChars) };
- jniContext->program = clCreateProgramWithSource( jniContext->context, 1, &sourceChars, sourceSize, &status);
- jenv->ReleaseStringUTFChars(source, sourceChars);
- ASSERT_CL("clCreateProgramWithSource()");
-
- status = clBuildProgram(jniContext->program, jniContext->deviceIdc, jniContext->deviceIds, NULL, NULL, NULL);
+
+ jniContext->program = CLHelper::compile(jenv, jniContext->context, jniContext->deviceIdc, jniContext->deviceIds, source, NULL, &status);
if(status == CL_BUILD_PROGRAM_FAILURE) {
- cl_int logStatus;
- size_t buildLogSize = 0;
- status = clGetProgramBuildInfo(jniContext->program, jniContext->deviceIds[0],
- CL_PROGRAM_BUILD_LOG, buildLogSize, NULL, &buildLogSize);
- ASSERT_CL("clGetProgramBuildInfo()");
- char * buildLog = new char[buildLogSize];
- CHECK(buildLog == NULL, "Failed to allocate host memory. (buildLog)");
- memset(buildLog, 0, buildLogSize);
- status = clGetProgramBuildInfo (jniContext->program, jniContext->deviceIds[0],
- CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL);
- ASSERT_CL("clGetProgramBuildInfo()");
-
- fprintf(stderr, "clBuildProgram failed");
- fprintf(stderr, "\n************************************************\n");
- fprintf(stderr, "%s", buildLog);
- fprintf(stderr, "\n************************************************\n\n\n");
- delete []buildLog;
return(0);
}
@@ -1557,7 +1374,7 @@ JNIEXPORT jlong JNICALL Java_com_amd_aparapi_KernelRunner_buildProgramJNI(JNIEnv
}
jniContext->commandQueues= new cl_command_queue[jniContext->deviceIdc];
- for (int dev=0; dev < jniContext->deviceIdc; dev++){
+ for (unsigned dev=0; dev < jniContext->deviceIdc; dev++){
jniContext->commandQueues[dev]=clCreateCommandQueue(jniContext->context, (cl_device_id)jniContext->deviceIds[dev],
queue_props,
&status);
@@ -1613,7 +1430,8 @@ JNIEXPORT jlong JNICALL Java_com_amd_aparapi_KernelRunner_buildProgramJNI(JNIEnv
// this is called once when the arg list is first determined for this kernel
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_setArgsJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jobjectArray argArray, jint argc) {
+JNI_JAVA(jint, KernelRunner, setArgsJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jobjectArray argArray, jint argc) {
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
cl_int status = CL_SUCCESS;
if (jniContext != NULL){
@@ -1723,18 +1541,13 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_setArgsJNI(JNIEnv *jenv
-JNIEXPORT jstring JNICALL Java_com_amd_aparapi_KernelRunner_getExtensionsJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
+JNI_JAVA(jstring, KernelRunner, getExtensionsJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
jstring jextensions = NULL;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext != NULL){
- size_t retvalsize = 0;
cl_int status = CL_SUCCESS;
- status = clGetDeviceInfo(jniContext->deviceIds[0], CL_DEVICE_EXTENSIONS, 0, NULL, &retvalsize);
- ASSERT_CL("clGetDeviceInfo()");
- char* extensions = new char[retvalsize];
- clGetDeviceInfo(jniContext->deviceIds[0], CL_DEVICE_EXTENSIONS, retvalsize, extensions, NULL);
- jextensions = jenv->NewStringUTF(extensions);
- delete [] extensions;
+ jextensions = CLHelper::getExtensions(jenv, jniContext->deviceIds[0], &status);
}
return jextensions;
}
@@ -1763,7 +1576,8 @@ KernelArg* getArgForBuffer(JNIEnv* jenv, JNIContext* jniContext, jobject buffer)
}
// Called as a result of Kernel.get(someArray)
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jobject buffer) {
+JNI_JAVA(jint, KernelRunner, getJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jobject buffer) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext != NULL){
@@ -1809,63 +1623,9 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getJNI(JNIEnv *jenv, jo
return 0;
}
-jobject createInstance(JNIEnv *jenv, char* className, char *signature, ... ){
- jclass theClass = jenv->FindClass(className);
- if (theClass == NULL || jenv->ExceptionCheck()) {
- jenv->ExceptionDescribe();
- jenv->ExceptionClear();
- fprintf(stderr, "bummer! getting '%s'\n", className);
- return(NULL);
- }
-
- jmethodID constructor= jenv->GetMethodID(theClass,"<init>",signature);
- if (constructor == NULL || jenv->ExceptionCheck()) {
- jenv->ExceptionDescribe();
- jenv->ExceptionClear();
- fprintf(stderr, "bummer getting constructor from '%s' with signature! '%s' \n", className, signature);
- return(NULL);
- }
- va_list argp;
- va_start(argp, signature);
- jobject instance = jenv->NewObjectV(theClass, constructor, argp);
- if (instance == NULL || jenv->ExceptionCheck()) {
- jenv->ExceptionDescribe();
- jenv->ExceptionClear();
- fprintf(stderr, "bummer invoking constructor from '%s' with signature! '%s' \n", className, signature);
- }
- va_end(argp);
- return(instance);
-}
-
-void callVoid(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...){
- jclass theClass = jenv->GetObjectClass(instance);
- if (theClass == NULL || jenv->ExceptionCheck()) {
- jenv->ExceptionDescribe();
- jenv->ExceptionClear();
- fprintf(stderr, "bummer! getting class from instance\n");
- return;
- }
- jmethodID methodId= jenv->GetMethodID(theClass,methodName,methodSignature);
- if (methodId == NULL || jenv->ExceptionCheck()) {
- jenv->ExceptionDescribe();
- jenv->ExceptionClear();
- fprintf(stderr, "bummer getting method '%s %s' from instance \n", methodName, methodSignature);
- return;
- }
- va_list argp;
- va_start(argp, methodSignature);
- jenv->CallVoidMethodV(instance, methodId, argp);
- if (jenv->ExceptionCheck()) {
- jenv->ExceptionDescribe(); /* write to console */
- jenv->ExceptionClear();
- fprintf(stderr, "bummer calling '%s %s'\n", methodName, methodSignature);
- }
- va_end(argp);
- return;
-}
jobject createProfileInfo(JNIEnv *jenv, ProfileInfo &profileInfo){
- jobject profileInstance = createInstance(jenv, "com/amd/aparapi/ProfileInfo", "(Ljava/lang/String;IJJJJ)V",
+ jobject profileInstance = JNIHelper::createInstance(jenv, "com/amd/aparapi/ProfileInfo", "(Ljava/lang/String;IJJJJ)V",
((jstring)(profileInfo.name==NULL?NULL:jenv->NewStringUTF(profileInfo.name))),
((jint)profileInfo.type),
((jlong)profileInfo.start),
@@ -1875,12 +1635,13 @@ jobject createProfileInfo(JNIEnv *jenv, ProfileInfo &profileInfo){
return(profileInstance);
}
-JNIEXPORT jobject JNICALL Java_com_amd_aparapi_KernelRunner_getProfileInfoJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
+JNI_JAVA(jobject, KernelRunner, getProfileInfoJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
jobject returnList = NULL;
if (jniContext != NULL){
- returnList = createInstance(jenv, "java/util/ArrayList", "()V");
+ returnList = JNIHelper::createInstance(jenv, "java/util/ArrayList", "()V");
if (jniContext->isProfilingEnabled()){
for (jint i=0; i<jniContext->argc; i++){
@@ -1888,14 +1649,14 @@ JNIEXPORT jobject JNICALL Java_com_amd_aparapi_KernelRunner_getProfileInfoJNI(JN
if (arg->isArray()){
if (arg->isMutableByKernel() && arg->value.ref.write.valid){
jobject writeProfileInfo = createProfileInfo(jenv, arg->value.ref.write);
- callVoid(jenv, returnList, "add", "(Ljava/lang/Object;)Z", writeProfileInfo);
+ JNIHelper::callVoid(jenv, returnList, "add", "(Ljava/lang/Object;)Z", writeProfileInfo);
}
}
}
for (jint pass=0; pass<jniContext->passes; pass++){
jobject executeProfileInfo = createProfileInfo(jenv, jniContext->exec[pass]);
- callVoid(jenv, returnList, "add", "(Ljava/lang/Object;)Z", executeProfileInfo);
+ JNIHelper::callVoid(jenv, returnList, "add", "(Ljava/lang/Object;)Z", executeProfileInfo);
}
for (jint i=0; i<jniContext->argc; i++){
@@ -1903,7 +1664,7 @@ JNIEXPORT jobject JNICALL Java_com_amd_aparapi_KernelRunner_getProfileInfoJNI(JN
if (arg->isArray()){
if (arg->isReadByKernel() && arg->value.ref.read.valid){
jobject readProfileInfo = createProfileInfo(jenv, arg->value.ref.read);
- callVoid(jenv, returnList, "add", "(Ljava/lang/Object;)Z", readProfileInfo);
+ JNIHelper::callVoid(jenv, returnList, "add", "(Ljava/lang/Object;)Z", readProfileInfo);
}
}
}
@@ -1913,7 +1674,8 @@ JNIEXPORT jobject JNICALL Java_com_amd_aparapi_KernelRunner_getProfileInfoJNI(JN
}
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxComputeUnitsJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
+JNI_JAVA(jint, KernelRunner, getMaxComputeUnitsJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext != NULL){
@@ -1923,7 +1685,8 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxComputeUnitsJNI(J
}
}
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxWorkItemDimensionsJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
+JNI_JAVA(jint, KernelRunner, getMaxWorkItemDimensionsJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext != NULL){
@@ -1933,7 +1696,8 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxWorkItemDimension
}
}
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxWorkGroupSizeJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
+JNI_JAVA(jint, KernelRunner, getMaxWorkGroupSizeJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
if (jniContext != NULL){
@@ -1943,10 +1707,11 @@ JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxWorkGroupSizeJNI(
}
}
-JNIEXPORT jint JNICALL Java_com_amd_aparapi_KernelRunner_getMaxWorkItemSizeJNI(JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jint _index) {
+JNI_JAVA(jint, KernelRunner, getMaxWorkItemSizeJNI)
+ (JNIEnv *jenv, jobject jobj, jlong jniContextHandle, jint _index) {
cl_int status = CL_SUCCESS;
JNIContext* jniContext = JNIContext::getJNIContext(jniContextHandle);
- if (jniContext != NULL && _index >=0 && _index <= jniContext->maxWorkItemDimensions){
+ if (jniContext != NULL && _index >=0 && _index <= (int)(jniContext->maxWorkItemDimensions)){
return(jniContext->maxWorkItemSizes[_index]);
}else{
return(0);
diff --git a/com.amd.aparapi.jni/src/cpp/aparapi.h b/com.amd.aparapi.jni/src/cpp/aparapi.h
new file mode 100644
index 0000000000000000000000000000000000000000..53be6ed76cd560f0960777ac9c1226eac6ad1f86
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/aparapi.h
@@ -0,0 +1,42 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#ifndef APARAPI_H
+#define APARAPI_H
+
+#endif // APARAPI_H
diff --git a/com.amd.aparapi.jni/src/cpp/clHelper.cpp b/com.amd.aparapi.jni/src/cpp/clHelper.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..d7c5ccaa9df8859cf86b7ec48d794fe6f64a3cc3
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/clHelper.cpp
@@ -0,0 +1,162 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#include "common.h"
+
+#define CLHELPER_SOURCE
+#include "clHelper.h"
+
+const char *CLHelper::errString(cl_int status){
+ static struct { cl_int code; const char *msg; } error_table[] = {
+ { CL_SUCCESS, "success" },
+ { CL_DEVICE_NOT_FOUND, "device not found", },
+ { CL_DEVICE_NOT_AVAILABLE, "device not available", },
+ { CL_COMPILER_NOT_AVAILABLE, "compiler not available", },
+ { CL_MEM_OBJECT_ALLOCATION_FAILURE, "mem object allocation failure", },
+ { CL_OUT_OF_RESOURCES, "out of resources", },
+ { CL_OUT_OF_HOST_MEMORY, "out of host memory", },
+ { CL_PROFILING_INFO_NOT_AVAILABLE, "profiling not available", },
+ { CL_MEM_COPY_OVERLAP, "memcopy overlaps", },
+ { CL_IMAGE_FORMAT_MISMATCH, "image format mismatch", },
+ { CL_IMAGE_FORMAT_NOT_SUPPORTED, "image format not supported", },
+ { CL_BUILD_PROGRAM_FAILURE, "build program failed", },
+ { CL_MAP_FAILURE, "map failed", },
+ { CL_INVALID_VALUE, "invalid value", },
+ { CL_INVALID_DEVICE_TYPE, "invalid device type", },
+ { CL_INVALID_PLATFORM, "invlaid platform",},
+ { CL_INVALID_DEVICE, "invalid device",},
+ { CL_INVALID_CONTEXT, "invalid context",},
+ { CL_INVALID_QUEUE_PROPERTIES, "invalid queue properties",},
+ { CL_INVALID_COMMAND_QUEUE, "invalid command queue",},
+ { CL_INVALID_HOST_PTR, "invalid host ptr",},
+ { CL_INVALID_MEM_OBJECT, "invalid mem object",},
+ { CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, "invalid image format descriptor ",},
+ { CL_INVALID_IMAGE_SIZE, "invalid image size",},
+ { CL_INVALID_SAMPLER, "invalid sampler",},
+ { CL_INVALID_BINARY, "invalid binary",},
+ { CL_INVALID_BUILD_OPTIONS, "invalid build options",},
+ { CL_INVALID_PROGRAM, "invalid program ",},
+ { CL_INVALID_PROGRAM_EXECUTABLE, "invalid program executable",},
+ { CL_INVALID_KERNEL_NAME, "invalid kernel name",},
+ { CL_INVALID_KERNEL_DEFINITION, "invalid definition",},
+ { CL_INVALID_KERNEL, "invalid kernel",},
+ { CL_INVALID_ARG_INDEX, "invalid arg index",},
+ { CL_INVALID_ARG_VALUE, "invalid arg value",},
+ { CL_INVALID_ARG_SIZE, "invalid arg size",},
+ { CL_INVALID_KERNEL_ARGS, "invalid kernel args",},
+ { CL_INVALID_WORK_DIMENSION , "invalid work dimension",},
+ { CL_INVALID_WORK_GROUP_SIZE, "invalid work group size",},
+ { CL_INVALID_WORK_ITEM_SIZE, "invalid work item size",},
+ { CL_INVALID_GLOBAL_OFFSET, "invalid global offset",},
+ { CL_INVALID_EVENT_WAIT_LIST, "invalid event wait list",},
+ { CL_INVALID_EVENT, "invalid event",},
+ { CL_INVALID_OPERATION, "invalid operation",},
+ { CL_INVALID_GL_OBJECT, "invalid gl object",},
+ { CL_INVALID_BUFFER_SIZE, "invalid buffer size",},
+ { CL_INVALID_MIP_LEVEL, "invalid mip level",},
+ { CL_INVALID_GLOBAL_WORK_SIZE, "invalid global work size",},
+ { 0, NULL },
+ };
+ static char unknown[25];
+ int ii;
+
+ for (ii = 0; error_table[ii].msg != NULL; ii++) {
+ if (error_table[ii].code == status) {
+ return error_table[ii].msg;
+ }
+ }
+#ifdef _WIN32
+ _snprintf(unknown, sizeof unknown, "unknown error %d", status);
+#else
+ snprintf(unknown, sizeof(unknown), "unknown error %d", status);
+#endif
+ return unknown;
+}
+
+void CLHelper::getBuildErr(JNIEnv *jenv, cl_device_id deviceId, cl_program program, jstring *log){
+ size_t buildLogSize = 0;
+ clGetProgramBuildInfo(program, deviceId, CL_PROGRAM_BUILD_LOG, buildLogSize, NULL, &buildLogSize);
+ char * buildLog = new char[buildLogSize];
+ memset(buildLog, 0, buildLogSize);
+ clGetProgramBuildInfo (program, deviceId, CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL);
+ fprintf(stderr, "clBuildProgram failed");
+ fprintf(stderr, "\n************************************************\n");
+ fprintf(stderr, "%s", buildLog);
+ fprintf(stderr, "\n************************************************\n\n\n");
+ if (log != NULL){
+ *log = jenv->NewStringUTF(buildLog);
+ }
+ delete []buildLog;
+}
+
+cl_program CLHelper::compile(JNIEnv *jenv, cl_context context, size_t deviceCount, cl_device_id* deviceIds, jstring source, jstring* log, cl_int* status){
+ const char *sourceChars = jenv->GetStringUTFChars(source, NULL);
+ size_t sourceSize[] = { strlen(sourceChars) };
+ cl_program program = clCreateProgramWithSource(context, 1, &sourceChars, sourceSize, status);
+ jenv->ReleaseStringUTFChars(source, sourceChars);
+ *status = clBuildProgram(program, deviceCount, deviceIds, NULL, NULL, NULL);
+ if(*status == CL_BUILD_PROGRAM_FAILURE) {
+ getBuildErr(jenv, *deviceIds, program, log);
+ }
+ return(program);
+}
+
+jstring CLHelper::getExtensions(JNIEnv *jenv, cl_device_id deviceId, cl_int *status){
+ jstring jextensions = NULL;
+ size_t retvalsize = 0;
+ *status = clGetDeviceInfo(deviceId, CL_DEVICE_EXTENSIONS, 0, NULL, &retvalsize);
+ if (*status == CL_SUCCESS){
+ char* extensions = new char[retvalsize];
+ *status = clGetDeviceInfo(deviceId, CL_DEVICE_EXTENSIONS, retvalsize, extensions, NULL);
+ if (*status == CL_SUCCESS){
+ jextensions = jenv->NewStringUTF(extensions);
+ }
+ delete [] extensions;
+ }
+ return jextensions;
+}
+
+/*
+Program::Program(cl_context context, cl_device_id deviceId, cl_command_queue commandQueue, cl_program program):
+ context(context),
+ deviceId(deviceId),
+ commandQueue(commandQueue),
+ program(program){
+}
+*/
+
diff --git a/com.amd.aparapi.jni/src/cpp/clHelper.h b/com.amd.aparapi.jni/src/cpp/clHelper.h
new file mode 100644
index 0000000000000000000000000000000000000000..d4ec57ab7923690a50c222b6cc1a9d934bce879e
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/clHelper.h
@@ -0,0 +1,72 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#ifndef CLHELPER_H
+#define CLHELPER_H
+/*
+class Program{
+ private:
+ cl_context context;
+ cl_device_id deviceId;
+ cl_command_queue commandQueue;
+ cl_program program;
+ public:
+ Program(cl_context context, cl_device_id deviceId, cl_command_queue commandQueue, cl_program program);
+};
+
+class Buf{
+ private:
+ cl_mem mem;
+ void *ptr;
+ size_t size;
+ cl_uint mask;
+ cl_uint bits;
+};
+*/
+
+
+class CLHelper{
+ public:
+ static const char *errString(cl_int status);
+ static void getBuildErr(JNIEnv *jenv, cl_device_id deviceId, cl_program program, jstring *log);
+ static cl_program compile(JNIEnv *jenv, cl_context context, size_t deviceCount, cl_device_id* deviceId, jstring source, jstring* log, cl_int *status);
+ static jstring getExtensions(JNIEnv *jenv, cl_device_id deviceId, cl_int *status);
+};
+
+#endif // CLHELPER_H
+
diff --git a/com.amd.aparapi.jni/src/cpp/common.h b/com.amd.aparapi.jni/src/cpp/common.h
new file mode 100644
index 0000000000000000000000000000000000000000..0ee588e6121a147cfa771ddc0e55db6ea2811768
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/common.h
@@ -0,0 +1,73 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#ifndef COMMON_H
+#define COMMON_H
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <time.h>
+
+#ifndef __APPLE__
+#include <malloc.h>
+#endif
+
+#include <sys/types.h>
+#ifndef _WIN32
+#include <unistd.h>
+#endif
+
+#ifndef __APPLE__
+#include <CL/cl.h>
+#else
+#include <cl.h>
+#endif
+
+#include <jni.h>
+
+#if defined (_WIN32)
+#include "windows.h"
+#define alignedMalloc(size, alignment)\
+ _aligned_malloc(size, alignment)
+#else
+#define alignedMalloc(size, alignment)\
+ memalign(alignment, size)
+#endif
+
+#endif // COMMON_H
diff --git a/com.amd.aparapi.jni/src/cpp/jniHelper.cpp b/com.amd.aparapi.jni/src/cpp/jniHelper.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..3997bd5b08c3ccce7317ddb7c3aaec9bf1c4d20f
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/jniHelper.cpp
@@ -0,0 +1,445 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#include "common.h"
+#define JNI_SOURCE
+#include "jniHelper.h"
+
+void JNIHelper::callVoid(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return;
+ }
+ jmethodID methodId= jenv->GetMethodID(theClass,methodName,methodSignature);
+ if (methodId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting method '%s','%s' from instance \n", methodName, methodSignature);
+ return;
+ }
+ va_list argp;
+ va_start(argp, methodSignature);
+ jenv->CallVoidMethodV(instance, methodId, argp);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe(); /* write to console */
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer calling '%s %s'\n", methodName, methodSignature);
+ }
+ va_end(argp);
+ return;
+}
+
+jobject JNIHelper::callObject(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return NULL;
+ }
+ jmethodID methodId= jenv->GetMethodID(theClass,methodName,methodSignature);
+ if (methodId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting method '%s','%s' from instance \n", methodName, methodSignature);
+ return NULL;
+ }
+ va_list argp;
+ va_start(argp, methodSignature);
+ jobject value = jenv->CallObjectMethodV(instance, methodId, argp);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe(); /* write to console */
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer calling '%s %s'\n", methodName, methodSignature);
+ return NULL;
+ }
+ va_end(argp);
+ return value;
+}
+
+jlong JNIHelper::callLong(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return 0L;
+ }
+ jmethodID methodId= jenv->GetMethodID(theClass,methodName,methodSignature);
+ if (methodId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting method '%s','%s' from instance \n", methodName, methodSignature);
+ return 0L;
+ }
+ va_list argp;
+ va_start(argp, methodSignature);
+ jlong value = jenv->CallLongMethodV(instance, methodId, argp);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe(); /* write to console */
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer calling '%s %s'\n", methodName, methodSignature);
+ return 0L;
+ }
+ va_end(argp);
+ return value;
+}
+
+void JNIHelper::setInstanceFieldInt(JNIEnv* jenv, jobject instance, char *fieldName, jint value){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"I");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field '%s' \n", fieldName);
+ return;
+ }
+ jenv->SetIntField(instance, fieldId, value);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer setting int field '%s' \n", fieldName);
+ return;
+ }
+}
+
+void JNIHelper::setInstanceFieldLong(JNIEnv* jenv, jobject instance, char *fieldName, jlong value){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"J");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting long field '%s' \n", fieldName);
+ return;
+ }
+ jenv->SetLongField(instance, fieldId, value);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer setting long field '%s' \n", fieldName);
+ return;
+ }
+}
+void JNIHelper::setInstanceFieldBoolean(JNIEnv* jenv, jobject instance, char *fieldName, jboolean value){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"Z");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting boolean field id '%s' \n", fieldName);
+ return;
+ }
+ jenv->SetBooleanField(instance, fieldId, value);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer setting boolean field '%s' \n", fieldName);
+ return;
+ }
+}
+
+void JNIHelper::setInstanceFieldObject(JNIEnv* jenv, jobject instance, char *fieldName, char *signature, jobject value){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName, signature);
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting %s object '%s' \n", signature, fieldName);
+ return;
+ }
+ jenv->SetObjectField(instance, fieldId, value);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer setting %s object '%s' \n", signature, fieldName);
+ return;
+ }
+}
+
+jobject JNIHelper::getStaticFieldObject(JNIEnv *jenv, char *className, char *fieldName, char *signature){
+ jclass theClass = jenv->FindClass(className);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting '%s'\n", className);
+ return(NULL);
+ }
+ jfieldID fieldId= jenv->GetStaticFieldID(theClass,fieldName,signature);
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting static field '%s' from '%s' with signature! '%s' \n", fieldName, className, signature);
+ return(NULL);
+ }
+
+ jobject value = jenv->GetStaticObjectField(NULL, fieldId);
+ if (value == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting static field value '%s' from '%s' with signature! '%s' \n", fieldName, className, signature);
+ return(NULL);
+ }
+
+ return(value);
+}
+jobject JNIHelper::createInstance(JNIEnv *jenv, char* className, char *signature, ... ){
+ jclass theClass = jenv->FindClass(className);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting '%s'\n", className);
+ return(NULL);
+ }
+
+ jmethodID constructor= jenv->GetMethodID(theClass,"<init>",signature);
+ if (constructor == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting constructor from '%s' with signature! '%s' \n", className, signature);
+ return(NULL);
+ }
+ va_list argp;
+ va_start(argp, signature);
+ jobject instance = jenv->NewObjectV(theClass, constructor, argp);
+ if (instance == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer invoking constructor from '%s' with signature! '%s' \n", className, signature);
+ }
+ va_end(argp);
+ return(instance);
+}
+
+jint JNIHelper::getInstanceFieldInt(JNIEnv *jenv, jobject instance, char *fieldName){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return 0;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"I");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field '%s' \n", fieldName);
+ return 0;
+ }
+ jint value= jenv->GetIntField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field value '%s' \n", fieldName);
+ return 0;
+ }
+ return(value);
+}
+
+jfloat JNIHelper::getInstanceFieldFloat(JNIEnv *jenv, jobject instance, char *fieldName){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return 0;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"F");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field '%s' \n", fieldName);
+ return 0;
+ }
+ jfloat value= jenv->GetFloatField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field value '%s' \n", fieldName);
+ return 0;
+ }
+ return(value);
+}
+
+jdouble JNIHelper::getInstanceFieldDouble(JNIEnv *jenv, jobject instance, char *fieldName){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return 0;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"D");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field '%s' \n", fieldName);
+ return 0;
+ }
+ jdouble value= jenv->GetDoubleField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field value '%s' \n", fieldName);
+ return 0;
+ }
+ return(value);
+}
+jshort JNIHelper::getInstanceFieldShort(JNIEnv *jenv, jobject instance, char *fieldName){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return 0;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"H");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field '%s' \n", fieldName);
+ return 0;
+ }
+ jshort value= jenv->GetShortField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting int field value '%s' \n", fieldName);
+ return 0;
+ }
+ return(value);
+}
+
+jlong JNIHelper::getInstanceFieldLong(JNIEnv *jenv, jobject instance, char *fieldName){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return 0;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName,"J");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting long field '%s' \n", fieldName);
+ return 0;
+ }
+ jlong value= jenv->GetLongField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting long field value '%s' \n", fieldName);
+ return 0;
+ }
+ return(value);
+}
+jobject JNIHelper::getInstanceFieldObject(JNIEnv *jenv, jobject instance, char *fieldName, char *signature){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return NULL;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName, signature);
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting object field '%s' \n", fieldName);
+ return NULL;
+ }
+ jobject value= jenv->GetObjectField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting object field '%s' \n", fieldName);
+ return NULL;
+ }
+ return(value);
+}
+jboolean JNIHelper::getInstanceFieldBoolean(JNIEnv *jenv, jobject instance, char *fieldName){
+ jclass theClass = jenv->GetObjectClass(instance);
+ if (theClass == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer! getting class from instance\n");
+ return NULL;
+ }
+ jfieldID fieldId= jenv->GetFieldID(theClass,fieldName, "Z");
+ if (fieldId == NULL || jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting boolean field '%s' \n", fieldName);
+ return NULL;
+ }
+ jboolean value= jenv->GetBooleanField(instance, fieldId);
+ if (jenv->ExceptionCheck()) {
+ jenv->ExceptionDescribe();
+ jenv->ExceptionClear();
+ fprintf(stderr, "bummer getting boolean field '%s' \n", fieldName);
+ return NULL;
+ }
+ return(value);
+}
diff --git a/com.amd.aparapi.jni/src/cpp/jniHelper.h b/com.amd.aparapi.jni/src/cpp/jniHelper.h
new file mode 100644
index 0000000000000000000000000000000000000000..eb34a5da3117bea176b09f3e7282babf18d1de8c
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/jniHelper.h
@@ -0,0 +1,103 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#ifndef JNIHELPER_H
+#define JNIHELPER_H
+
+#include <jni.h>
+
+class JNIHelper{
+ public:
+ static void callVoid(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...);
+ static jlong callLong(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...);
+ static jobject callObject(JNIEnv *jenv, jobject instance, char *methodName, char *methodSignature, ...);
+
+ static jobject JNIHelper::createInstance(JNIEnv *jenv, char *className, char *signature, ... );
+
+ static jobject getStaticFieldObject(JNIEnv *jenv, char *className, char *fieldName, char *signature);
+
+ static jint getInstanceFieldInt(JNIEnv *jenv, jobject instance, char *fieldName);
+ static jfloat getInstanceFieldFloat(JNIEnv *jenv, jobject instance, char *fieldName);
+ static jdouble getInstanceFieldDouble(JNIEnv *jenv, jobject instance, char *fieldName);
+ static jshort getInstanceFieldShort(JNIEnv *jenv, jobject instance, char *fieldName);
+ static jlong getInstanceFieldLong(JNIEnv *jenv, jobject instance, char *fieldName);
+ static jboolean getInstanceFieldBoolean(JNIEnv *jenv, jobject instance, char *fieldName);
+ static jobject getInstanceFieldObject(JNIEnv *jenv, jobject instance, char *fieldName, char *signature);
+
+ static void setInstanceFieldInt(JNIEnv* jenv, jobject instance, char *fieldName, jint value);
+ static void setInstanceFieldLong(JNIEnv* jenv, jobject instance, char *fieldName, jlong value);
+ static void setInstanceFieldBoolean(JNIEnv* jenv, jobject instance, char *fieldName, jboolean value);
+ static void setInstanceFieldObject(JNIEnv* jenv, jobject instance, char *fieldName, char *signature, jobject value);
+};
+
+#define JNIExceptionChecker(){\
+ fprintf(stderr, "line %d\n", __LINE__);\
+ if ((jenv)->ExceptionOccurred()) {\
+ (jenv)->ExceptionDescribe(); /* write to console */\
+ (jenv)->ExceptionClear();\
+ }\
+}
+
+#define CHECK_NO_RETURN(condition, msg) if(condition){\
+ fprintf(stderr, "!!!!!!! %s failed !!!!!!!\n", msg);\
+}
+
+#define CHECK_RETURN(condition, msg, val) if(condition){\
+ fprintf(stderr, "!!!!!!! %s failed !!!!!!!\n", msg);\
+ return val;\
+}
+
+#define CHECK(condition, msg) CHECK_RETURN(condition, msg, 0)
+
+#define ASSERT_CL_NO_RETURN(msg) if (status != CL_SUCCESS){\
+ fprintf(stderr, "!!!!!!! %s failed: %s\n", msg, CLHelper::errString(status));\
+}
+
+#define ASSERT_CL_RETURN(msg, val) if (status != CL_SUCCESS){\
+ ASSERT_CL_NO_RETURN(msg)\
+ return val;\
+}
+
+#define ASSERT_CL(msg) ASSERT_CL_RETURN(msg, 0)
+
+#define PRINT_CL_ERR(status, msg) fprintf(stderr, "!!!!!!! %s failed %s\n", msg, CLHelper::errString(status));
+
+#define ASSERT_FIELD(id) CHECK_NO_RETURN(id##FieldID == 0, "No such field as " #id)
+
+#endif // JNIHELPER_H
+
diff --git a/com.amd.aparapi.jni/src/cpp/opencljni.cpp b/com.amd.aparapi.jni/src/cpp/opencljni.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..e7ba1c279877ac92d396877de0e4614ee46f0cd8
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/opencljni.cpp
@@ -0,0 +1,715 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#include "common.h"
+#include "jniHelper.h"
+#include "clHelper.h"
+
+#define OPENCLJNI_SOURCE
+#include "opencljni.h"
+
+#include "com_amd_aparapi_OpenCLJNI.h"
+
+#define JNI_JAVA(type, className, methodName) JNIEXPORT type JNICALL Java_com_amd_aparapi_##className##_##methodName
+#define isset(bits, token) (((bits) & com_amd_aparapi_OpenCLJNI_##token##_BIT) ==com_amd_aparapi_OpenCLJNI_##token##_BIT)
+#define set(bits, token) (bits) |= com_amd_aparapi_OpenCLJNI_##token##_BIT
+#define reset(bits, token) (bits) &= ~com_amd_aparapi_OpenCLJNI_##token##_BIT
+
+class OpenCLBits{
+ public:
+ static void describeBits(JNIEnv *jenv, jlong bits){
+ fprintf(stderr, " %lx ", bits);
+ if (isset(bits, READONLY)){
+ fprintf(stderr, "readonly ");
+ }
+ if (isset(bits, WRITEONLY)){
+ fprintf(stderr, "writeonly ");
+ }
+ if (isset(bits, READWRITE)){
+ fprintf(stderr, "readwrite ");
+ }
+ if (isset(bits, ARRAY)){
+ fprintf(stderr, "array ");
+ }
+ if (isset(bits, PRIMITIVE)){
+ fprintf(stderr, "primitive ");
+ }
+ if (isset(bits, FLOAT)){
+ fprintf(stderr, "float ");
+ }
+ if (isset(bits, SHORT)){
+ fprintf(stderr, "short ");
+ }
+ if (isset(bits, LONG)){
+ fprintf(stderr, "long ");
+ }
+ if (isset(bits, DOUBLE)){
+ fprintf(stderr, "double ");
+ }
+ if (isset(bits, INT)){
+ fprintf(stderr, "int ");
+ }
+ if (isset(bits, MEM_COPY)){
+ fprintf(stderr, "copy ");
+ }
+ if (isset(bits, GLOBAL)){
+ fprintf(stderr, "global ");
+ }
+ if (isset(bits, LOCAL)){
+ fprintf(stderr, "local ");
+ }
+ if (isset(bits, MEM_DIRTY)){
+ fprintf(stderr, "dirty ");
+ }
+ if (isset(bits, MEM_ENQUEUED)){
+ fprintf(stderr, "enqueued ");
+ }
+ if (isset(bits, ARG)){
+ fprintf(stderr, "arg ");
+ }
+ }
+};
+
+class OpenCLDevice{
+ public:
+ static jobject getPlatformInstance(JNIEnv *jenv, jobject deviceInstance){
+ return(JNIHelper::getInstanceFieldObject(jenv, deviceInstance, "platform", "Lcom/amd/aparapi/OpenCLPlatform;"));
+ }
+ static cl_device_id getDeviceId(JNIEnv *jenv, jobject deviceInstance){
+ return((cl_device_id)JNIHelper::getInstanceFieldLong(jenv, deviceInstance, "deviceId"));
+ }
+};
+class OpenCLPlatform{
+ public:
+ static cl_platform_id getPlatformId(JNIEnv *jenv, jobject platformInstance){
+ return((cl_platform_id)JNIHelper::getInstanceFieldLong(jenv, platformInstance, "platformId"));
+ }
+};
+class OpenCLProgram{
+ public:
+ static jobject create(JNIEnv *jenv, cl_program program, cl_command_queue queue, cl_context context, jobject deviceInstance, jstring source, jstring log){
+ return(JNIHelper::createInstance(jenv, "com/amd/aparapi/OpenCLProgram", "(JJJLcom/amd/aparapi/OpenCLDevice;Ljava/lang/String;Ljava/lang/String;)V", (jlong)program, (jlong)queue, (jlong)context, deviceInstance, source, log));
+ }
+ static cl_context getContext(JNIEnv *jenv, jobject programInstance){
+ return((cl_context) JNIHelper::getInstanceFieldLong(jenv, programInstance, "contextId"));
+ }
+ static cl_program getProgram(JNIEnv *jenv, jobject programInstance){
+ return((cl_program) JNIHelper::getInstanceFieldLong(jenv, programInstance, "programId"));
+ }
+ static cl_command_queue getCommandQueue(JNIEnv *jenv, jobject programInstance){
+ return((cl_command_queue)JNIHelper::getInstanceFieldLong(jenv, programInstance, "queueId"));
+ }
+};
+class OpenCLKernel{
+ public:
+ static jobject create(JNIEnv *jenv, cl_kernel kernel, jobject programInstance, jstring name, jobject args){
+ return(JNIHelper::createInstance(jenv, "com/amd/aparapi/OpenCLKernel", "(JLcom/amd/aparapi/OpenCLProgram;Ljava/lang/String;Ljava/util/List;)V", (jlong)kernel, programInstance, name, args));
+ }
+
+ static cl_kernel getKernel(JNIEnv *jenv, jobject kernelInstance){
+ return((cl_kernel) JNIHelper::getInstanceFieldLong(jenv, kernelInstance, "kernelId"));
+ }
+
+ static jobject getProgramInstance(JNIEnv *jenv, jobject kernelInstance){
+ return(JNIHelper::getInstanceFieldObject(jenv, kernelInstance, "program", "Lcom/amd/aparapi/OpenCLProgram;"));
+ }
+
+ static jobjectArray getArgsArray(JNIEnv *jenv, jobject kernelInstance){
+ return(reinterpret_cast<jobjectArray> (JNIHelper::getInstanceFieldObject(jenv, kernelInstance, "args", "[Lcom/amd/aparapi/OpenCLArg;")));
+ }
+};
+
+class OpenCLMem{
+ public:
+ static jobject create(JNIEnv *jenv){
+ return(JNIHelper::createInstance(jenv, "Lcom/amd/aparapi/OpenCLMem;", "()V"));
+ }
+ static cl_uint bitsToOpenCLMask(jlong argBits ){
+ cl_uint mask = CL_MEM_USE_HOST_PTR;
+ if (isset(argBits, READONLY)){
+ mask|= CL_MEM_READ_ONLY;
+ }else if (isset(argBits, READWRITE)){
+ mask|= CL_MEM_READ_WRITE;
+ } else if (isset(argBits, WRITEONLY)){
+ mask|= CL_MEM_WRITE_ONLY;
+ }
+ return(mask);
+ }
+
+ static jsize getArraySizeInBytes(JNIEnv *jenv, jarray array, jlong argBits){
+ jsize arrayLen = jenv->GetArrayLength(array);
+ jsize sizeInBytes = 0;
+ if (isset(argBits, FLOAT) || isset(argBits, INT)){
+ sizeInBytes = arrayLen *4;
+ }else if (isset(argBits, DOUBLE) || isset(argBits, LONG)){
+ sizeInBytes = arrayLen *8;
+ }else if (isset(argBits, SHORT)){
+ sizeInBytes = arrayLen *2;
+ }
+ fprintf(stderr, "size of new Mem arg is %d\n", sizeInBytes);
+ return(sizeInBytes);
+ }
+ static void *pin(JNIEnv *jenv, jarray array, jlong *memBits){
+ jboolean isCopy;
+ void *ptr = jenv->GetPrimitiveArrayCritical(array,&isCopy);
+ if (memBits != NULL){
+ if(0){OpenCLBits::describeBits(jenv, *memBits);fprintf(stderr, " before \n");}
+ if (isCopy){
+ set(*memBits, MEM_COPY);
+ }else{
+ reset(*memBits, MEM_COPY);
+ }
+ if(0){OpenCLBits::describeBits(jenv, *memBits);fprintf(stderr, " after \n");}
+ }
+ return(ptr);
+ }
+
+ static void unpin(JNIEnv *jenv, jarray array, void *ptr, jlong *memBits){
+ if (isset(*memBits, WRITEONLY)){
+ jenv->ReleasePrimitiveArrayCritical(array, ptr,JNI_ABORT);
+ }else{
+ jenv->ReleasePrimitiveArrayCritical(array, ptr, 0);
+ }
+ }
+
+
+ static jobject create(JNIEnv *jenv, cl_context context, jlong argBits, jarray array){
+ jsize sizeInBytes = getArraySizeInBytes(jenv, array, argBits);
+
+ cl_int status = CL_SUCCESS;
+ void *ptr = OpenCLMem::pin(jenv, array, &argBits);
+ cl_mem mem=clCreateBuffer(context, bitsToOpenCLMask(argBits), sizeInBytes, ptr, &status);
+ if (status != CL_SUCCESS){
+ fprintf(stderr, "buffer creation failed!\n");
+ }
+
+ jobject memInstance = OpenCLMem::create(jenv);
+ fprintf(stderr, "created a new mem object!\n");
+ OpenCLMem::setAddress(jenv, memInstance, ptr);
+ OpenCLMem::setInstance(jenv, memInstance, array);
+ OpenCLMem::setSizeInBytes(jenv, memInstance, sizeInBytes);
+ OpenCLMem::setBits(jenv, memInstance, argBits);
+ OpenCLMem::setMem(jenv, memInstance, argBits, mem);
+ fprintf(stderr, "initiated mem object!\n");
+ return(memInstance);
+ }
+ static jlong getBits(JNIEnv *jenv, jobject memInstance){
+ return(JNIHelper::getInstanceFieldLong(jenv, memInstance, "bits"));
+ }
+ static void setBits(JNIEnv *jenv, jobject memInstance, jlong bits){
+ JNIHelper::setInstanceFieldLong(jenv, memInstance, "bits", bits);
+ }
+ static void *getAddress(JNIEnv *jenv, jobject memInstance){
+ return((void*)JNIHelper::getInstanceFieldLong(jenv, memInstance, "address"));
+ }
+ static void setAddress(JNIEnv *jenv, jobject memInstance, void *address){
+ JNIHelper::setInstanceFieldLong(jenv, memInstance, "address", (jlong)address);
+ }
+ static void setInstance(JNIEnv *jenv, jobject memInstance, jobject instance){
+ JNIHelper::setInstanceFieldObject(jenv, memInstance, "instance", "Ljava/lang/Object;", instance);
+ }
+ static void setSizeInBytes(JNIEnv *jenv, jobject memInstance, jint sizeInBytes){
+ JNIHelper::setInstanceFieldInt(jenv, memInstance, "sizeInBytes", sizeInBytes);
+ }
+ static size_t getSizeInBytes(JNIEnv *jenv, jobject memInstance){
+ return((size_t)JNIHelper::getInstanceFieldInt(jenv, memInstance, "sizeInBytes"));
+ }
+ static jobject getInstance(JNIEnv *jenv, jobject memInstance){
+ return(JNIHelper::getInstanceFieldObject(jenv, memInstance, "instance", "Ljava/lang/Object;"));
+ }
+
+ static cl_mem getMem(JNIEnv *jenv, jobject memInstance, jlong argBits){
+ cl_mem mem = 0;
+
+ if (isset(argBits, READONLY)){
+ mem = (cl_mem)JNIHelper::getInstanceFieldLong(jenv, memInstance, "readOnlyMemId");
+ } else if (isset(argBits, READWRITE)){
+ mem = (cl_mem)JNIHelper::getInstanceFieldLong(jenv, memInstance, "readWriteMemId");
+ } else if (isset(argBits, WRITEONLY)){
+ mem = (cl_mem)JNIHelper::getInstanceFieldLong(jenv, memInstance, "writeOnlyMemId");
+ }
+ return(mem);
+ }
+
+ static void setMem(JNIEnv *jenv, jobject memInstance, jlong argBits, cl_mem mem){
+ if (isset(argBits, READONLY)){
+ JNIHelper::setInstanceFieldLong(jenv, memInstance, "readOnlyMemId", (jlong)mem);
+ } else if (isset(argBits, READWRITE)){
+ JNIHelper::setInstanceFieldLong(jenv, memInstance, "readWriteMemId", (jlong)mem);
+ } else if (isset(argBits, WRITEONLY)){
+ JNIHelper::setInstanceFieldLong(jenv, memInstance, "writeOnlyMemId", (jlong)mem);
+ }
+ }
+
+ static void describe(JNIEnv *jenv, jobject memInstance){
+ jlong memBits = OpenCLMem::getBits(jenv, memInstance);
+ OpenCLBits::describeBits(jenv, memBits);
+ fprintf(stderr, "\n");
+ }
+
+};
+
+class OpenCLArg{
+ public:
+ static jlong getBits(JNIEnv *jenv, jobject argInstance){
+ return(JNIHelper::getInstanceFieldLong(jenv, argInstance, "bits"));
+ }
+ static void setBits(JNIEnv *jenv, jobject argInstance, jlong bits){
+ JNIHelper::setInstanceFieldLong(jenv, argInstance, "bits", bits);
+ }
+ static jobject getMemInstance(JNIEnv *jenv, jobject argInstance){
+ return(JNIHelper::getInstanceFieldObject(jenv, argInstance, "memVal", "Lcom/amd/aparapi/OpenCLMem;"));
+ }
+ static void setMemInstance(JNIEnv *jenv, jobject argInstance, jobject memInstance){
+ JNIHelper::setInstanceFieldObject(jenv, argInstance, "memVal", "Lcom/amd/aparapi/OpenCLMem;",memInstance);
+ }
+ static void describe(JNIEnv *jenv, jobject argDef, jint argIndex){
+ jlong argBits = OpenCLArg::getBits(jenv, argDef);
+ fprintf(stderr, " %d ", argIndex);
+ OpenCLBits::describeBits(jenv, argBits);
+ fprintf(stderr, "\n");
+ }
+};
+class OpenCLRange{
+ public:
+ static jint getDims(JNIEnv *jenv, jobject rangeInstance){
+ return(JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "dims"));
+ }
+
+ static void fill(JNIEnv *jenv, jobject rangeInstance, jint dims, size_t* offsets, size_t* globalDims, size_t* localDims){
+ if (dims >0){
+ offsets[0]= 0;
+ localDims[0]=JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "localSize_0");
+ globalDims[0]=JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "globalSize_0");
+ if (dims >1){
+ offsets[1]= 0;
+ localDims[1]=JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "localSize_1");
+ globalDims[1]=JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "globalSize_1");
+ if (dims >2){
+ offsets[2]= 0;
+ localDims[2]=JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "localSize_2");
+ globalDims[2]=JNIHelper::getInstanceFieldInt(jenv, rangeInstance, "globalSize_2");
+ }
+ }
+ }
+ }
+};
+
+JNI_JAVA(jobject, OpenCLJNI, createProgram)
+ (JNIEnv *jenv, jobject jobj, jobject deviceInstance, jstring source) {
+
+ jobject platformInstance = OpenCLDevice::getPlatformInstance(jenv, deviceInstance);
+ cl_platform_id platformId = OpenCLPlatform::getPlatformId(jenv, platformInstance);
+ cl_device_id deviceId = OpenCLDevice::getDeviceId(jenv, deviceInstance);
+ cl_int status = CL_SUCCESS;
+ cl_device_type deviceType;
+ clGetDeviceInfo(deviceId, CL_DEVICE_TYPE, sizeof(deviceType), &deviceType, NULL);
+ //fprintf(stderr, "device[%d] CL_DEVICE_TYPE = %x\n", deviceId, deviceType);
+
+
+ cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platformId, 0 };
+ cl_context_properties* cprops = (NULL == platformId) ? NULL : cps;
+ cl_context context = clCreateContextFromType( cprops, deviceType, NULL, NULL, &status);
+
+
+ jstring log=NULL;
+ cl_program program = CLHelper::compile(jenv, context, 1, &deviceId, source, &log, &status);
+ cl_command_queue queue = NULL;
+ if(status == CL_SUCCESS) {
+ cl_command_queue_properties queue_props = CL_QUEUE_PROFILING_ENABLE;
+ queue = clCreateCommandQueue(context, deviceId, queue_props, &status);
+ }else{
+ fprintf(stderr, "queue creation seems to have failed\n");
+
+ }
+ jobject programInstance = OpenCLProgram::create(jenv, program, queue, context, deviceInstance, source, log);
+
+ return(programInstance);
+ }
+
+JNI_JAVA(jobject, OpenCLJNI, createKernel)
+ (JNIEnv *jenv, jobject jobj, jobject programInstance, jstring name, jobject args) {
+ cl_context context = OpenCLProgram::getContext(jenv, programInstance);
+ cl_program program = OpenCLProgram::getProgram(jenv, programInstance);
+ cl_int status = CL_SUCCESS;
+ const char *nameChars = jenv->GetStringUTFChars(name, NULL);
+ fprintf(stderr, "tring to extract kernel '%s'\n", nameChars);
+ cl_kernel kernel = clCreateKernel(program, nameChars, &status);
+ jenv->ReleaseStringUTFChars(name, nameChars);
+
+ if (kernel == NULL){
+ fprintf(stderr, "kernel is null!\n");
+ }
+
+ jobject kernelInstance = NULL;
+ if (status == CL_SUCCESS){
+ kernelInstance = OpenCLKernel::create(jenv, kernel, programInstance, name, args);
+ }else{
+ fprintf(stderr, "kernel creation seems to have failed\n");
+ }
+ return(kernelInstance);
+
+ }
+
+
+
+void putArgs(JNIEnv *jenv, cl_context context, cl_kernel kernel, cl_command_queue commandQueue, cl_event *events, jint *eventc, jint argIndex, jobject argDef, jobject arg){
+ if(0){
+ fprintf(stderr, "putArgs ");
+ OpenCLArg::describe(jenv, argDef, argIndex);
+ }
+ cl_int status = CL_SUCCESS;
+ jlong argBits = OpenCLArg::getBits(jenv, argDef);
+ if (isset(argBits, ARRAY)){ // global check?
+ jobject memInstance = OpenCLArg::getMemInstance(jenv, argDef);
+ if (memInstance == NULL){
+ // first call?
+ memInstance = OpenCLMem::create(jenv, context, argBits, (jarray)arg);
+ OpenCLArg::setMemInstance(jenv, argDef, memInstance);
+ }else{
+ // check of argBits == memInstance.argBits
+ // we need to pin it
+ jboolean isCopy;
+ void *ptr = OpenCLMem::pin(jenv, (jarray)arg,&argBits);
+ void *oldPtr = OpenCLMem::getAddress(jenv, memInstance);
+ if (ptr !=oldPtr){
+ fprintf(stderr, "ptr moved from %lx to %lx\n", oldPtr, ptr);
+ cl_mem mem = OpenCLMem::getMem(jenv, memInstance, argBits);
+ status = clReleaseMemObject(mem);
+ memInstance = OpenCLMem::create(jenv, context, argBits, (jarray)arg);
+ OpenCLArg::setMemInstance(jenv, argDef, memInstance);
+ }
+ OpenCLArg::setBits(jenv, argDef, argBits);
+ }
+ cl_mem mem = OpenCLMem::getMem(jenv, memInstance, argBits);
+ cl_int status = CL_SUCCESS;
+ if (isset(argBits, READONLY)|isset(argBits, READWRITE)){ // kernel reads this so enqueue a write
+ void *ptr= OpenCLMem::getAddress(jenv, memInstance);
+ size_t sizeInBytes= OpenCLMem::getSizeInBytes(jenv, memInstance);
+ jlong memBits = OpenCLMem::getBits(jenv, memInstance);
+ set(memBits, MEM_ENQUEUED);
+ OpenCLMem::setBits(jenv, memInstance, memBits);
+ status = clEnqueueWriteBuffer(commandQueue, mem, CL_FALSE, 0, sizeInBytes, ptr, *eventc, (*eventc)==0?NULL:events, &events[*eventc]);
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error enqueuing write %s!\n", CLHelper::errString(status));
+ }else{
+ if(0)fprintf(stderr, "enqueued write eventc = %d!\n", *eventc);
+ }
+ (*eventc)++;
+ }
+ status = clSetKernelArg(kernel, argIndex, sizeof(cl_mem), (void *)&(mem));
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error setting arg %d %s!\n", argIndex, CLHelper::errString(status));
+ }else{
+ if(0)fprintf(stderr, "set arg = %d!\n", argIndex);
+ }
+ }else if (isset(argBits, PRIMITIVE)){
+ if (isset(argBits, INT)){
+ cl_int value = JNIHelper::getInstanceFieldInt(jenv, arg, "value");
+ status = clSetKernelArg(kernel, argIndex, sizeof(value), (void *)&(value));
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error setting int arg %d %d %s!\n", argIndex, value, CLHelper::errString(status));
+ }else{
+ if(0)fprintf(stderr, "set arg = %d to %d!\n", argIndex, value);
+ }
+ }else if (isset(argBits, FLOAT)){
+ cl_float value = JNIHelper::getInstanceFieldFloat(jenv, arg, "value");
+ status = clSetKernelArg(kernel, argIndex, sizeof(value), (void *)&(value));
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error setting int arg %d %f %s!\n", argIndex, value, CLHelper::errString(status));
+ }else{
+ if(0)fprintf(stderr, "set arg = %d to %f!\n", argIndex, value);
+ }
+
+ }else if (isset(argBits, DOUBLE)){
+ cl_double value = JNIHelper::getInstanceFieldDouble(jenv, arg, "value");
+ status = clSetKernelArg(kernel, argIndex, sizeof(value), (void *)&(value));
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error setting double arg %d %lf %s!\n", argIndex, value, CLHelper::errString(status));
+ }else{
+ if(0)fprintf(stderr, "set arg = %d to %lf!\n", argIndex, value);
+ }
+
+ }
+ }
+}
+
+
+
+
+void getArgs(JNIEnv *jenv, cl_context context, cl_command_queue commandQueue, cl_event *events, jint *eventc, jint argIndex, jobject argDef, jobject arg){
+ if (0){
+ fprintf(stderr, "post ");
+ OpenCLArg::describe(jenv, argDef, argIndex);
+ }
+ jlong argBits = OpenCLArg::getBits(jenv, argDef);
+ if (isset(argBits, ARRAY)){
+ jobject memInstance = OpenCLArg::getMemInstance(jenv, argDef);
+ if (memInstance == NULL){
+ fprintf(stderr, "mem instance not set\n");
+ }else{
+ if(0)fprintf(stderr, "retrieved mem instance\n");
+ }
+ void *ptr= OpenCLMem::getAddress(jenv, memInstance);
+ if (isset(argBits, WRITEONLY)|isset(argBits, READWRITE)){
+
+ cl_mem mem = OpenCLMem::getMem(jenv, memInstance, argBits);
+
+ size_t sizeInBytes= OpenCLMem::getSizeInBytes(jenv, memInstance);
+ if (0){
+ fprintf(stderr, "about to enqueu read eventc = %d!\n", *eventc);
+ }
+ cl_int status = clEnqueueReadBuffer(commandQueue, mem, CL_FALSE, 0, sizeInBytes, ptr ,*eventc, (*eventc)==0?NULL:events, &events[*eventc]);
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error enqueuing read %s!\n", CLHelper::errString(status));
+ }else{
+ if (0){
+ fprintf(stderr, "enqueued read eventc = %d!\n", *eventc);
+ }
+ }
+ (*eventc)++;
+ }
+
+ jobject arrayInstance = OpenCLMem::getInstance(jenv, memInstance);
+
+ OpenCLMem::unpin(jenv, (jarray)arrayInstance, ptr, &argBits);
+
+
+ reset(argBits, MEM_ENQUEUED); //<----- BAD
+ reset(argBits, MEM_COPY); //<----
+ OpenCLMem::setBits(jenv, memInstance, argBits); // WHAT? GRF
+ }
+}
+
+JNI_JAVA(void, OpenCLJNI, invoke)
+ (JNIEnv *jenv, jobject jobj, jobject kernelInstance, jobjectArray argArray) {
+
+
+ cl_kernel kernel = OpenCLKernel::getKernel(jenv, kernelInstance);
+ jobject programInstance = OpenCLKernel::getProgramInstance(jenv, kernelInstance);
+ jobjectArray argDefsArray = OpenCLKernel::getArgsArray(jenv, kernelInstance);
+
+ cl_context context =OpenCLProgram::getContext(jenv, programInstance);
+ cl_command_queue commandQueue = OpenCLProgram::getCommandQueue(jenv, programInstance);
+
+
+ // walk through the args creating buffers when needed
+ // we use the bitfields to determine which is which
+ // note that argArray[0] is the range then 1,2,3 etc matches argDefsArray[0,1,2]
+ jsize argc = jenv->GetArrayLength(argDefsArray);
+ if (0) fprintf(stderr, "argc = %d\n", argc);
+ jint reads=0;
+ jint writes=0;
+ for (jsize argIndex=0; argIndex<argc; argIndex++){
+ jobject argDef = jenv->GetObjectArrayElement(argDefsArray, argIndex);
+ jlong argBits = OpenCLArg::getBits(jenv, argDef);
+ if (isset(argBits, READONLY)){
+ reads++;
+ }
+ if (isset(argBits, READWRITE)){
+ reads++;
+ writes++;
+ }
+ if (isset(argBits, WRITEONLY)){
+ writes++;
+ }
+ }
+
+ if (0) fprintf(stderr, "reads=%d writes=%d\n", reads, writes);
+ cl_event * events= new cl_event[reads+writes+1];
+
+ jint eventc =0;
+
+ for (jsize argIndex=0; argIndex<argc; argIndex++){
+ jobject argDef = jenv->GetObjectArrayElement(argDefsArray, argIndex);
+ jobject arg = jenv->GetObjectArrayElement(argArray, argIndex+1);
+ putArgs(jenv, context, kernel, commandQueue, events, &eventc, argIndex, argDef, arg);
+ }
+
+ jobject rangeInstance = jenv->GetObjectArrayElement(argArray, 0);
+ jint dims = OpenCLRange::getDims(jenv, rangeInstance);
+
+ size_t *offsets = new size_t[dims];
+ size_t *globalDims = new size_t[dims];
+ size_t *localDims = new size_t[dims];
+ OpenCLRange::fill(jenv, rangeInstance, dims, offsets, globalDims, localDims);
+
+ cl_int status = CL_SUCCESS;
+
+ if(0) fprintf(stderr, "Exec %d\n", eventc);
+ status = clEnqueueNDRangeKernel(
+ commandQueue,
+ kernel,
+ dims,
+ offsets,
+ globalDims,
+ localDims,
+ eventc, // count Of events to wait for
+ eventc==0?NULL:events, // address of events to wait for
+ &events[eventc]);
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error enqueuing execute %s !\n", CLHelper::errString(status));
+ }else{
+ if (0) fprintf(stderr, "success enqueuing execute eventc= %d !\n", eventc);
+ }
+ eventc++;
+
+ for (jsize argIndex=0; argIndex<argc; argIndex++){
+ jobject argDef = jenv->GetObjectArrayElement(argDefsArray, argIndex);
+ jobject arg = jenv->GetObjectArrayElement(argArray, argIndex+1);
+ getArgs(jenv, context, commandQueue, events, &eventc, argIndex, argDef, arg);
+ }
+ status = clWaitForEvents(eventc, events);
+ if (status != CL_SUCCESS) {
+ fprintf(stderr, "error waiting for events !\n");
+ }
+ }
+
+JNI_JAVA(jobject, OpenCLJNI, getPlatforms)
+ (JNIEnv *jenv, jobject jobj) {
+ jobject platformListInstance = JNIHelper::createInstance(jenv, "java/util/ArrayList", "()V");
+ cl_int status = CL_SUCCESS;
+ cl_uint platformc;
+
+ status = clGetPlatformIDs(0, NULL, &platformc);
+ //fprintf(stderr, "There are %d platforms\n", platformc);
+ cl_platform_id* platformIds = new cl_platform_id[platformc];
+ status = clGetPlatformIDs(platformc, platformIds, NULL);
+
+ if (status == CL_SUCCESS){
+ for (unsigned platformIdx = 0; platformIdx < platformc; ++platformIdx) {
+ char platformVersionName[512];
+ status = clGetPlatformInfo(platformIds[platformIdx], CL_PLATFORM_VERSION, sizeof(platformVersionName), platformVersionName, NULL);
+ if ( !strncmp(platformVersionName, "OpenCL 1.1", 10)
+ || !strncmp(platformVersionName, "OpenCL 1.2", 10)) {
+ char platformVendorName[512];
+ status = clGetPlatformInfo(platformIds[platformIdx], CL_PLATFORM_VENDOR, sizeof(platformVendorName), platformVendorName, NULL);
+ //fprintf(stderr, "platform vendor %d %s\n", platformIdx, platformVendorName);
+ //fprintf(stderr, "platform version %d %s\n", platformIdx, platformVersionName);
+ jobject platformInstance = JNIHelper::createInstance(jenv, "com/amd/aparapi/OpenCLPlatform", "(JLjava/lang/String;Ljava/lang/String;)V",
+ (jlong)platformIds[platformIdx],
+ jenv->NewStringUTF(platformVersionName),
+ jenv->NewStringUTF(platformVendorName));
+ JNIHelper::callVoid(jenv, platformListInstance, "add", "(Ljava/lang/Object;)Z", platformInstance);
+
+ cl_uint deviceIdc;
+ cl_device_type requestedDeviceType =CL_DEVICE_TYPE_CPU |CL_DEVICE_TYPE_GPU ;
+ status = clGetDeviceIDs(platformIds[platformIdx], requestedDeviceType, 0, NULL, &deviceIdc);
+ if (status == CL_SUCCESS && deviceIdc >0 ){
+ cl_device_id* deviceIds = new cl_device_id[deviceIdc];
+ status = clGetDeviceIDs(platformIds[platformIdx], requestedDeviceType, deviceIdc, deviceIds, NULL);
+ if (status == CL_SUCCESS){
+ for (unsigned deviceIdx=0; deviceIdx<deviceIdc; deviceIdx++){
+
+ cl_device_type deviceType;
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_TYPE, sizeof(deviceType), &deviceType, NULL);
+ jobject deviceTypeEnumInstance = JNIHelper::getStaticFieldObject(jenv, "com/amd/aparapi/OpenCLDevice$TYPE", "UNKNOWN", "Lcom/amd/aparapi/OpenCLDevice$TYPE;");
+ //fprintf(stderr, "device[%d] CL_DEVICE_TYPE = ", deviceIdx);
+ if (deviceType & CL_DEVICE_TYPE_DEFAULT) {
+ deviceType &= ~CL_DEVICE_TYPE_DEFAULT;
+ // fprintf(stderr, "Default ");
+ }
+ if (deviceType & CL_DEVICE_TYPE_CPU) {
+ deviceType &= ~CL_DEVICE_TYPE_CPU;
+ //fprintf(stderr, "CPU ");
+ deviceTypeEnumInstance = JNIHelper::getStaticFieldObject(jenv, "com/amd/aparapi/OpenCLDevice$TYPE", "CPU", "Lcom/amd/aparapi/OpenCLDevice$TYPE;");
+ }
+ if (deviceType & CL_DEVICE_TYPE_GPU) {
+ deviceType &= ~CL_DEVICE_TYPE_GPU;
+ //fprintf(stderr, "GPU ");
+ deviceTypeEnumInstance = JNIHelper::getStaticFieldObject(jenv, "com/amd/aparapi/OpenCLDevice$TYPE", "GPU", "Lcom/amd/aparapi/OpenCLDevice$TYPE;");
+ }
+ if (deviceType & CL_DEVICE_TYPE_ACCELERATOR) {
+ deviceType &= ~CL_DEVICE_TYPE_ACCELERATOR;
+ //fprintf(stderr, "Accelerator ");
+ }
+ //fprintf(stderr, "(0x%llx) ", deviceType);
+ //fprintf(stderr, "\n");
+
+
+ jobject deviceInstance = JNIHelper::createInstance(jenv, "com/amd/aparapi/OpenCLDevice", "(Lcom/amd/aparapi/OpenCLPlatform;JLcom/amd/aparapi/OpenCLDevice$TYPE;)V",
+ platformInstance,
+ (jlong)deviceIds[deviceIdx],
+ deviceTypeEnumInstance);
+ JNIHelper::callVoid(jenv, platformInstance, "add", "(Lcom/amd/aparapi/OpenCLDevice;)V", deviceInstance);
+
+
+ cl_uint maxComputeUnits;
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(maxComputeUnits), &maxComputeUnits, NULL);
+ //fprintf(stderr, "device[%d] CL_DEVICE_MAX_COMPUTE_UNITS = %u\n", deviceIdx, maxComputeUnits);
+ JNIHelper::callVoid(jenv, deviceInstance, "setMaxComputeUnits", "(I)V", maxComputeUnits);
+
+
+
+ cl_uint maxWorkItemDimensions;
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(maxWorkItemDimensions), &maxWorkItemDimensions, NULL);
+ //fprintf(stderr, "device[%d] CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS = %u\n", deviceIdx, maxWorkItemDimensions);
+ JNIHelper::callVoid(jenv, deviceInstance, "setMaxWorkItemDimensions", "(I)V", maxWorkItemDimensions);
+
+ size_t *maxWorkItemSizes = new size_t[maxWorkItemDimensions];
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*maxWorkItemDimensions, maxWorkItemSizes, NULL);
+ for (unsigned dimIdx=0; dimIdx<maxWorkItemDimensions; dimIdx++){
+ //fprintf(stderr, "device[%d] dim[%d] = %d\n", deviceIdx, dimIdx, maxWorkItemSizes[dimIdx]);
+ JNIHelper::callVoid(jenv, deviceInstance, "setMaxWorkItemSize", "(II)V", dimIdx,maxWorkItemSizes[dimIdx]);
+ }
+
+ size_t maxWorkGroupSize;
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGroupSize), &maxWorkGroupSize, NULL);
+ //fprintf(stderr, "device[%d] CL_DEVICE_MAX_GROUP_SIZE = %u\n", deviceIdx, maxWorkGroupSize);
+ JNIHelper::callVoid(jenv, deviceInstance, "setMaxWorkGroupSize", "(I)V", maxWorkGroupSize);
+
+ cl_ulong globalMemSize;
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(globalMemSize), &globalMemSize, NULL);
+ //fprintf(stderr, "device[%d] CL_DEVICE_GLOBAL_MEM_SIZE = %lu\n", deviceIdx, globalMemSize);
+ JNIHelper::callVoid(jenv, deviceInstance, "setGlobalMemSize", "(J)V", globalMemSize);
+
+ cl_ulong localMemSize;
+ status = clGetDeviceInfo(deviceIds[deviceIdx], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(localMemSize), &localMemSize, NULL);
+ //fprintf(stderr, "device[%d] CL_DEVICE_LOCAL_MEM_SIZE = %lu\n", deviceIdx, localMemSize);
+ JNIHelper::callVoid(jenv, deviceInstance, "setLocalMemSize", "(J)V", localMemSize);
+ }
+
+ }
+ }
+ }
+
+ }
+ }
+ return (platformListInstance);
+ }
+
diff --git a/com.amd.aparapi.jni/src/cpp/opencljni.h b/com.amd.aparapi.jni/src/cpp/opencljni.h
new file mode 100644
index 0000000000000000000000000000000000000000..7c12ebdfa4cf91ac6d1162f44f9c64232f34e2b9
--- /dev/null
+++ b/com.amd.aparapi.jni/src/cpp/opencljni.h
@@ -0,0 +1,42 @@
+/*
+ Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+ All rights reserved.
+
+ Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+ following conditions are met:
+
+ Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+ disclaimer.
+
+ Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+ disclaimer in the documentation and/or other materials provided with the distribution.
+
+ Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+ derived from this software without specific prior written permission.
+
+ THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+ INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+ SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+ SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+ If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+ laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730
+ through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of
+ the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce
+ Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export
+ Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in
+ Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2)
+ export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced
+ direct product is subject to national security controls as identified on the Commerce Control List (currently
+ found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional
+ information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry
+ and Security’s website at http://www.bis.doc.gov/.
+ */
+
+#ifndef OPENCLJNI_H
+#define OPENCLJNI_H
+
+#endif // OPENCLJNI_H
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java b/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java
index 8d6b45b43174b897ec2b5029b1a93a68bd4c4ddf..c9b8bf3ed7e3521004ae962847ca06c6eeff727a 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/ClassModel.java
@@ -471,8 +471,8 @@ class ClassModel{
}
}
if (inMethod) {
- methodDescription = new MethodDescription(className, methodName, stringStack.toArray(new String[0])[0], methodStack
- .toArray(new String[0]));
+ methodDescription = new MethodDescription(className, methodName, stringStack.toArray(new String[0])[0],
+ methodStack.toArray(new String[0]));
} else {
System.out.println("can't convert to a description");
}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java b/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java
index 38bdbdf101588b961451ca1ddba493305812f5da..e9b7608a49a7bcc5e9f319d6283d2087077107c0 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/Entrypoint.java
@@ -346,8 +346,8 @@ class Entrypoint{
ArrayList<FieldEntry> structMemberSet = superCandidate.getStructMembers();
for (FieldEntry f : structMemberSet) {
if (f.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(accessedFieldName)
- && f.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8().equals(
- field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8())) {
+ && f.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8()
+ .equals(field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8())) {
if (logger.isLoggable(Level.FINE)) {
logger.fine("Found match: " + accessedFieldName + " class: " + field.getClassEntry().getNameUTF8Entry().getUTF8()
@@ -377,8 +377,8 @@ class Entrypoint{
ArrayList<FieldEntry> structMemberSet = memberClassModel.getStructMembers();
for (FieldEntry f : structMemberSet) {
if (f.getNameAndTypeEntry().getNameUTF8Entry().getUTF8().equals(accessedFieldName)
- && f.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8().equals(
- field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8())) {
+ && f.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8()
+ .equals(field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8())) {
found = true;
}
}
@@ -467,8 +467,8 @@ class Entrypoint{
// Look for a intra-object call in a object member
if (m == null) {
for (ClassModel c : allFieldsClasses.values()) {
- if (c.getClassWeAreModelling().getName().equals(
- methodEntry.getClassEntry().getNameUTF8Entry().getUTF8().replace('/', '.'))) {
+ if (c.getClassWeAreModelling().getName()
+ .equals(methodEntry.getClassEntry().getNameUTF8Entry().getUTF8().replace('/', '.'))) {
m = c.getMethod(methodEntry, (methodCall instanceof I_INVOKESPECIAL) ? true : false);
assert m != null;
break;
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/InstructionSet.java b/com.amd.aparapi/src/java/com/amd/aparapi/InstructionSet.java
index 9502576ea4774f8f8667add9639d07e6aa7759a5..80ad5b31fee06b8bd9d56562841b59436b7a8069 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/InstructionSet.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/InstructionSet.java
@@ -40,9 +40,9 @@ package com.amd.aparapi;
import java.lang.reflect.Constructor;
import java.lang.reflect.InvocationTargetException;
-import com.amd.aparapi.ClassModel.ConstantPool;
import com.amd.aparapi.ClassModel.AttributePool.LocalVariableTableEntry;
import com.amd.aparapi.ClassModel.AttributePool.LocalVariableTableEntry.LocalVariableInfo;
+import com.amd.aparapi.ClassModel.ConstantPool;
import com.amd.aparapi.ClassModel.ConstantPool.Entry;
import com.amd.aparapi.ClassModel.ConstantPool.FieldEntry;
import com.amd.aparapi.ClassModel.ConstantPool.MethodEntry;
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/Kernel.java b/com.amd.aparapi/src/java/com/amd/aparapi/Kernel.java
index 865bdbe82c248125e510ccb688e7094feda989f4..ee94ea1e87ae86083b4977082f92c22aa5ac87e0 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/Kernel.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/Kernel.java
@@ -327,10 +327,9 @@ public abstract class Kernel implements Cloneable{
}
} catch (UnsatisfiedLinkError e) {
- logger
- .warning("Check your environment. Failed to load aparapi native library "
- + libName
- + " or possibly failed to locate opencl native library (opencl.dll/opencl.so). Ensure that both are in your PATH (windows) or in LD_LIBRARY_PATH (linux).");
+ logger.warning("Check your environment. Failed to load aparapi native library "
+ + libName
+ + " or possibly failed to locate opencl native library (opencl.dll/opencl.so). Ensure that both are in your PATH (windows) or in LD_LIBRARY_PATH (linux).");
openCLAvailable = false;
}
@@ -1991,11 +1990,11 @@ public abstract class Kernel implements Cloneable{
return (this);
}
- /**
- * Tag this array so that it is explicitly enqueued before the kernel is executed
- * @param array
- * @return This kernel so that we can use the 'fluent' style API
- */
+ /**
+ * Tag this array so that it is explicitly enqueued before the kernel is executed
+ * @param array
+ * @return This kernel so that we can use the 'fluent' style API
+ */
public Kernel put(char[] array) {
if (kernelRunner == null) {
kernelRunner = new KernelRunner(this);
@@ -2004,7 +2003,6 @@ public abstract class Kernel implements Cloneable{
return (this);
}
-
/**
* Enqueue a request to return this buffer from the GPU. This method blocks until the array is available.
* @param array
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java b/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java
index de4f6359d4f76515593ae204b7743302b5dc5272..e11f1ff1bb0bae90fb4dcba429d3acccbcfc138e 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/KernelWriter.java
@@ -43,11 +43,11 @@ import java.util.Iterator;
import java.util.List;
import java.util.Map;
-import com.amd.aparapi.ClassModel.ClassModelField;
import com.amd.aparapi.ClassModel.AttributePool.LocalVariableTableEntry;
-import com.amd.aparapi.ClassModel.AttributePool.RuntimeAnnotationsEntry;
import com.amd.aparapi.ClassModel.AttributePool.LocalVariableTableEntry.LocalVariableInfo;
+import com.amd.aparapi.ClassModel.AttributePool.RuntimeAnnotationsEntry;
import com.amd.aparapi.ClassModel.AttributePool.RuntimeAnnotationsEntry.AnnotationInfo;
+import com.amd.aparapi.ClassModel.ClassModelField;
import com.amd.aparapi.ClassModel.ConstantPool.FieldEntry;
import com.amd.aparapi.ClassModel.ConstantPool.MethodEntry;
import com.amd.aparapi.InstructionSet.AccessArrayElement;
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java b/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java
index 96106eea8247bef026ce48cfbdb4bc9a8c863c46..f2894849d81ce560870e4aa031935e73435d4f65 100644
--- a/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/MethodModel.java
@@ -44,15 +44,15 @@ import java.util.Iterator;
import java.util.LinkedList;
import java.util.List;
import java.util.Map;
-import java.util.Set;
import java.util.Map.Entry;
+import java.util.Set;
import java.util.logging.Level;
import java.util.logging.Logger;
-import com.amd.aparapi.ClassModel.ClassModelMethod;
-import com.amd.aparapi.ClassModel.ConstantPool;
import com.amd.aparapi.ClassModel.AttributePool.LocalVariableTableEntry;
import com.amd.aparapi.ClassModel.AttributePool.LocalVariableTableEntry.LocalVariableInfo;
+import com.amd.aparapi.ClassModel.ClassModelMethod;
+import com.amd.aparapi.ClassModel.ConstantPool;
import com.amd.aparapi.ClassModel.ConstantPool.FieldEntry;
import com.amd.aparapi.ClassModel.ConstantPool.MethodReferenceEntry;
import com.amd.aparapi.ClassModel.ConstantPool.MethodReferenceEntry.Arg;
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCL.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCL.java
new file mode 100644
index 0000000000000000000000000000000000000000..a934514ccc1dbfc842cd8f181dbe6eeb57a105ce
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCL.java
@@ -0,0 +1,76 @@
+package com.amd.aparapi;
+
+import java.lang.annotation.ElementType;
+import java.lang.annotation.Retention;
+import java.lang.annotation.RetentionPolicy;
+import java.lang.annotation.Target;
+
+public interface OpenCL<T> {
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface Put {
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface Get {
+ }
+
+ public @Target(ElementType.TYPE) @Retention(RetentionPolicy.RUNTIME) @interface Source {
+ String value();
+
+ }
+
+ public @Target(ElementType.TYPE) @Retention(RetentionPolicy.RUNTIME) @interface Resource {
+ String value();
+ }
+
+ public @Target(ElementType.METHOD) @Retention(RetentionPolicy.RUNTIME) @interface Kernel {
+ String value();
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface Arg {
+ String value();
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface GlobalReadWrite {
+ String value();
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface GlobalReadOnly {
+ String value();
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface GlobalWriteOnly {
+ String value();
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface Local {
+ String value();
+ }
+
+ public @Target(ElementType.PARAMETER) @Retention(RetentionPolicy.RUNTIME) @interface Constant {
+ String value();
+ }
+
+ public T put(float[] array);
+
+ public T put(int[] array);
+
+ public T put(short[] array);
+
+ public T put(char[] array);
+
+ public T put(boolean[] array);
+
+ public T put(double[] array);
+
+ public T get(float[] array);
+
+ public T get(int[] array);
+
+ public T get(short[] array);
+
+ public T get(char[] array);
+
+ public T get(boolean[] array);
+
+ public T get(double[] array);
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLAdapter.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLAdapter.java
new file mode 100644
index 0000000000000000000000000000000000000000..6e685b39c9a3794c8ea1ca08596ac12e67bb5a81
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLAdapter.java
@@ -0,0 +1,52 @@
+package com.amd.aparapi;
+
+public class OpenCLAdapter<T> implements OpenCL<T>{
+
+ public T put(float[] array) {
+ return ((T) this);
+ }
+
+ public T put(int[] array) {
+ return ((T) this);
+ }
+
+ public T put(short[] array) {
+ return ((T) this);
+ }
+
+ public T put(char[] array) {
+ return ((T) this);
+ }
+
+ public T put(boolean[] array) {
+ return ((T) this);
+ }
+
+ public T put(double[] array) {
+ return ((T) this);
+ }
+
+ public T get(float[] array) {
+ return ((T) this);
+ }
+
+ public T get(int[] array) {
+ return ((T) this);
+ }
+
+ public T get(short[] array) {
+ return ((T) this);
+ }
+
+ public T get(char[] array) {
+ return ((T) this);
+ }
+
+ public T get(boolean[] array) {
+ return ((T) this);
+ }
+
+ public T get(double[] array) {
+ return ((T) this);
+ }
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLArg.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLArg.java
new file mode 100644
index 0000000000000000000000000000000000000000..a48ec17a1e39dda92be865af9f049f952271ecb9
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLArg.java
@@ -0,0 +1,61 @@
+/**
+ *
+ */
+package com.amd.aparapi;
+
+public class OpenCLArg{
+ public OpenCLMem memVal;
+
+ private String name;
+
+ public long bits;
+
+ public OpenCLKernel kernel;
+
+ public OpenCLArg(String _name, long _bits) {
+
+ name = _name;
+ bits = _bits;
+ }
+
+ public String toString() {
+ StringBuilder argBuilder = new StringBuilder();
+ if ((bits & OpenCLJNI.GLOBAL_BIT) == OpenCLJNI.GLOBAL_BIT) {
+ argBuilder.append("__global ");
+ } else if ((bits & OpenCLJNI.LOCAL_BIT) == OpenCLJNI.LOCAL_BIT) {
+ argBuilder.append("__local ");
+ } else if ((bits & OpenCLJNI.CONST_BIT) == OpenCLJNI.CONST_BIT) {
+ argBuilder.append("__constant ");
+ } else if ((bits & OpenCLJNI.ARG_BIT) == OpenCLJNI.ARG_BIT) {
+ //
+ } else {
+ argBuilder.append("WHATISTHIS?");
+ }
+
+ if ((bits & OpenCLJNI.FLOAT_BIT) == OpenCLJNI.FLOAT_BIT) {
+ argBuilder.append("float ");
+ } else if ((bits & OpenCLJNI.INT_BIT) == OpenCLJNI.INT_BIT) {
+ argBuilder.append("int ");
+ } else if ((bits & OpenCLJNI.SHORT_BIT) == OpenCLJNI.SHORT_BIT) {
+ argBuilder.append("short ");
+ } else if ((bits & OpenCLJNI.DOUBLE_BIT) == OpenCLJNI.DOUBLE_BIT) {
+ argBuilder.append("double ");
+ } else if ((bits & OpenCLJNI.LONG_BIT) == OpenCLJNI.LONG_BIT) {
+ argBuilder.append("long ");
+ }
+
+ if ((bits & OpenCLJNI.ARRAY_BIT) == OpenCLJNI.ARRAY_BIT) {
+ argBuilder.append("*");
+ }
+ argBuilder.append(name);
+ if ((bits & OpenCLJNI.READONLY_BIT) == OpenCLJNI.READONLY_BIT) {
+ argBuilder.append(" /* readonly */");
+ } else if ((bits & OpenCLJNI.WRITEONLY_BIT) == OpenCLJNI.WRITEONLY_BIT) {
+ argBuilder.append(" /* writeonly */");
+ } else if ((bits & OpenCLJNI.READWRITE_BIT) == OpenCLJNI.READWRITE_BIT) {
+ argBuilder.append(" /* readwrite */");
+ }
+ return (argBuilder.toString());
+ }
+
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLDevice.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLDevice.java
new file mode 100644
index 0000000000000000000000000000000000000000..b2ee217ec981c7aa60511b15cc9f3692e570d78c
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLDevice.java
@@ -0,0 +1,436 @@
+package com.amd.aparapi;
+
+import java.io.BufferedReader;
+import java.io.IOException;
+import java.io.InputStream;
+import java.io.InputStreamReader;
+import java.lang.annotation.Annotation;
+import java.lang.reflect.InvocationHandler;
+import java.lang.reflect.Method;
+import java.lang.reflect.Proxy;
+import java.util.ArrayList;
+import java.util.HashMap;
+import java.util.List;
+import java.util.Map;
+
+public class OpenCLDevice{
+ static public enum TYPE {
+ UNKNOWN,
+ GPU,
+ CPU
+ };
+
+ private OpenCLPlatform platform;
+
+ private long deviceId;
+
+ private TYPE type = TYPE.UNKNOWN;
+
+ private int maxComputeUnits;
+
+ private int maxWorkItemDimensions;
+
+ private long localMemSize;
+
+ private long globalMemSize;
+
+ private int maxWorkGroupSize;
+
+ private int[] maxWorkItemSize = new int[] {
+ 0,
+ 0,
+ 0
+ };
+
+ OpenCLDevice(OpenCLPlatform _platform, long _deviceId, TYPE _type) {
+ platform = _platform;
+ deviceId = _deviceId;
+ type = _type;
+
+ }
+
+ public TYPE getType() {
+ return type;
+ }
+
+ public void setType(TYPE type) {
+ this.type = type;
+ }
+
+ public int getMaxComputeUnits() {
+ return maxComputeUnits;
+ }
+
+ public void setMaxComputeUnits(int _maxComputeUnits) {
+ maxComputeUnits = _maxComputeUnits;
+ }
+
+ public int getMaxWorkItemDimensions() {
+ return maxWorkItemDimensions;
+ }
+
+ public void setMaxWorkItemDimensions(int _maxWorkItemDimensions) {
+ maxWorkItemDimensions = _maxWorkItemDimensions;
+ }
+
+ public long getLocalMemSize() {
+ return localMemSize;
+ }
+
+ public void setLocalMemSize(long _localMemSize) {
+ localMemSize = _localMemSize;
+ }
+
+ public long getGlobalMemSize() {
+ return globalMemSize;
+ }
+
+ public void setGlobalMemSize(long _globalMemSize) {
+ globalMemSize = _globalMemSize;
+ }
+
+ public int getMaxWorkGroupSize() {
+ return maxWorkGroupSize;
+ }
+
+ public void setMaxWorkGroupSize(int _maxWorkGroupSize) {
+ maxWorkGroupSize = _maxWorkGroupSize;
+ }
+
+ public int[] getMaxWorkItemSize() {
+ return maxWorkItemSize;
+ }
+
+ public void setMaxWorkItemSize(int[] maxWorkItemSize) {
+ this.maxWorkItemSize = maxWorkItemSize;
+ }
+
+ public String toString() {
+ StringBuilder s = new StringBuilder("{");
+ boolean first = true;
+ for (int workItemSize : maxWorkItemSize) {
+ if (first) {
+ first = false;
+ } else {
+ s.append(", ");
+ }
+ s.append(workItemSize);
+ }
+ s.append("}");
+ return ("Device " + deviceId + "\n type:" + type + "\n maxComputeUnits=" + maxComputeUnits + "\n maxWorkItemDimensions="
+ + maxWorkItemDimensions + "\n maxWorkItemSizes=" + s + "\n maxWorkWorkGroupSize=" + maxWorkGroupSize
+ + "\n globalMemSize=" + globalMemSize + "\n localMemSize=" + localMemSize);
+ }
+
+ void setMaxWorkItemSize(int _dim, int _value) {
+ maxWorkItemSize[_dim] = _value;
+ }
+
+ public long getDeviceId() {
+ return (deviceId);
+ }
+
+ public OpenCLPlatform getPlatform() {
+ return (platform);
+ }
+
+ public static class OpenCLInvocationHandler<T extends OpenCL<T>> implements InvocationHandler{
+ private Map<String, OpenCLKernel> map;
+
+ private OpenCLProgram program;
+
+ public OpenCLInvocationHandler(OpenCLProgram _program, Map<String, OpenCLKernel> _map) {
+ program = _program;
+ map = _map;
+ }
+
+ @Override public Object invoke(Object proxy, Method method, Object[] args) throws Throwable {
+ OpenCLKernel kernel = map.get(method.getName());
+ if (kernel != null) {
+ // we have a kernel entrypoint bound
+ kernel.invoke(args);
+ } else if (method.getName().equals("put") || method.getName().equals("get")) {
+ for (Object arg : args) {
+ Class<?> argClass = arg.getClass();
+ if (argClass.isArray()) {
+ if (argClass.getComponentType().isPrimitive()) {
+ OpenCLMem mem = program.getMem(arg, 0L);
+ if (mem == null) {
+ throw new IllegalStateException("can't put/get an array that has never been passed to a kernel " + argClass);
+
+ }
+ if (method.getName().equals("put")) {
+ mem.bits |= OpenCLJNI.MEM_DIRTY_BIT;
+ } else {
+ OpenCLJNI.getJNI().getMem(program, mem);
+ }
+
+ } else {
+ throw new IllegalStateException("Only array args (of primitives) expected for put/get, cant deal with "
+ + argClass);
+
+ }
+ } else {
+ throw new IllegalStateException("Only array args expected for put/get, cant deal with " + argClass);
+ }
+ }
+ } else {
+ throw new IllegalStateException("How did we get here with method " + method.getName());
+
+ }
+ return proxy;
+ }
+
+ }
+
+ public List<OpenCLArg> getArgs(Method m) {
+ List<OpenCLArg> args = new ArrayList<OpenCLArg>();
+ Annotation[][] parameterAnnotations = m.getParameterAnnotations();
+ Class<?>[] parameterTypes = m.getParameterTypes();
+
+ for (int arg = 0; arg < parameterTypes.length; arg++) {
+ if (parameterTypes[arg].isAssignableFrom(Range.class)) {
+
+ } else {
+
+ long bits = 0L;
+ String name = null;
+ for (Annotation pa : parameterAnnotations[arg]) {
+ if (pa instanceof OpenCL.GlobalReadOnly) {
+ name = ((OpenCL.GlobalReadOnly) pa).value();
+ bits |= OpenCLJNI.GLOBAL_BIT | OpenCLJNI.READONLY_BIT;
+ } else if (pa instanceof OpenCL.GlobalWriteOnly) {
+ name = ((OpenCL.GlobalWriteOnly) pa).value();
+ bits |= OpenCLJNI.GLOBAL_BIT | OpenCLJNI.WRITEONLY_BIT;
+ } else if (pa instanceof OpenCL.GlobalReadWrite) {
+ name = ((OpenCL.GlobalReadWrite) pa).value();
+ bits |= OpenCLJNI.GLOBAL_BIT | OpenCLJNI.READWRITE_BIT;
+ } else if (pa instanceof OpenCL.Local) {
+ name = ((OpenCL.Local) pa).value();
+ bits |= OpenCLJNI.LOCAL_BIT;
+ } else if (pa instanceof OpenCL.Constant) {
+ name = ((OpenCL.Constant) pa).value();
+ bits |= OpenCLJNI.CONST_BIT | OpenCLJNI.READONLY_BIT;
+ } else if (pa instanceof OpenCL.Arg) {
+ name = ((OpenCL.Arg) pa).value();
+ bits |= OpenCLJNI.ARG_BIT;
+ }
+
+ }
+ if (parameterTypes[arg].isArray()) {
+ if (parameterTypes[arg].isAssignableFrom(float[].class)) {
+ bits |= OpenCLJNI.FLOAT_BIT | OpenCLJNI.ARRAY_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(int[].class)) {
+ bits |= OpenCLJNI.INT_BIT | OpenCLJNI.ARRAY_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(double[].class)) {
+ bits |= OpenCLJNI.DOUBLE_BIT | OpenCLJNI.ARRAY_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(short[].class)) {
+ bits |= OpenCLJNI.SHORT_BIT | OpenCLJNI.ARRAY_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(long[].class)) {
+ bits |= OpenCLJNI.LONG_BIT | OpenCLJNI.ARRAY_BIT;
+ }
+ } else if (parameterTypes[arg].isPrimitive()) {
+ if (parameterTypes[arg].isAssignableFrom(float.class)) {
+ bits |= OpenCLJNI.FLOAT_BIT | OpenCLJNI.PRIMITIVE_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(int.class)) {
+ bits |= OpenCLJNI.INT_BIT | OpenCLJNI.PRIMITIVE_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(double.class)) {
+ bits |= OpenCLJNI.DOUBLE_BIT | OpenCLJNI.PRIMITIVE_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(short.class)) {
+ bits |= OpenCLJNI.SHORT_BIT | OpenCLJNI.PRIMITIVE_BIT;
+ } else if (parameterTypes[arg].isAssignableFrom(long.class)) {
+ bits |= OpenCLJNI.LONG_BIT | OpenCLJNI.PRIMITIVE_BIT;
+ }
+ } else {
+ System.out.println("OUch!");
+ }
+ if (name == null) {
+ throw new IllegalStateException("no name!");
+ }
+ OpenCLArg kernelArg = new OpenCLArg(name, bits);
+ args.add(kernelArg);
+
+ }
+ }
+
+ return (args);
+ }
+
+ public <T extends OpenCL<T>> T create(Class<T> _interface) {
+
+ StringBuilder sourceBuilder = new StringBuilder();
+ Map<String, List<OpenCLArg>> kernelNameToArgsMap = new HashMap<String, List<OpenCLArg>>();
+ boolean interfaceIsAnnotated = false;
+ for (Annotation a : _interface.getAnnotations()) {
+ if (a instanceof OpenCL.Source) {
+ OpenCL.Source source = (OpenCL.Source) a;
+ sourceBuilder.append(source.value()).append("\n");
+ interfaceIsAnnotated = true;
+ } else if (a instanceof OpenCL.Resource) {
+ OpenCL.Resource sourceResource = (OpenCL.Resource) a;
+ InputStream stream = _interface.getClassLoader().getResourceAsStream(sourceResource.value());
+ if (stream != null) {
+
+ BufferedReader reader = new BufferedReader(new InputStreamReader(stream));
+
+ try {
+ for (String line = reader.readLine(); line != null; line = reader.readLine()) {
+ sourceBuilder.append(line).append("\n");
+ }
+ } catch (IOException e) {
+ // TODO Auto-generated catch block
+ e.printStackTrace();
+ }
+
+ }
+ interfaceIsAnnotated = true;
+ }
+ }
+
+ if (interfaceIsAnnotated) {
+ // just crawl the methods (non put or get) and create kernels
+ for (Method m : _interface.getDeclaredMethods()) {
+ if ((!m.getName().equals("put") && !m.getName().equals("get"))) {
+
+ List<OpenCLArg> args = getArgs(m);
+
+ kernelNameToArgsMap.put(m.getName(), args);
+ }
+ }
+ } else {
+ for (Method m : _interface.getDeclaredMethods()) {
+
+ for (Annotation a : m.getAnnotations()) {
+ // System.out.println(" annotation "+a);
+ // System.out.println(" annotation type " + a.annotationType());
+ if (a instanceof OpenCL.Kernel && (!m.getName().equals("put") && !m.getName().equals("get"))) {
+ sourceBuilder.append("__kernel void " + m.getName() + "(");
+ List<OpenCLArg> args = getArgs(m);
+
+ boolean first = true;
+ for (OpenCLArg arg : args) {
+ if (first) {
+ first = false;
+ } else {
+ sourceBuilder.append(",");
+ }
+ sourceBuilder.append("\n " + arg);
+ }
+
+ sourceBuilder.append(")");
+ OpenCL.Kernel kernel = (OpenCL.Kernel) a;
+ sourceBuilder.append(kernel.value());
+ kernelNameToArgsMap.put(m.getName(), args);
+ }
+ }
+ }
+ }
+
+ String source = sourceBuilder.toString();
+ System.out.println("opencl{\n" + source + "\n}opencl");
+
+ OpenCLProgram program = createProgram(source);
+
+ Map<String, OpenCLKernel> map = new HashMap<String, OpenCLKernel>();
+ for (String name : kernelNameToArgsMap.keySet()) {
+ OpenCLKernel kernel = program.createKernel(name, kernelNameToArgsMap.get(name));
+ if (kernel == null) {
+ throw new IllegalStateException("kernel is null");
+ }
+ map.put(name, kernel);
+ }
+
+ OpenCLInvocationHandler<T> invocationHandler = new OpenCLInvocationHandler<T>(program, map);
+ T instance = (T) Proxy.newProxyInstance(OpenCLDevice.class.getClassLoader(), new Class[] {
+ _interface,
+ OpenCL.class
+ }, invocationHandler);
+ return instance;
+
+ }
+
+ interface DeviceFilter{
+ boolean match(OpenCLDevice _device);
+ }
+
+ interface DeviceComparitor{
+ OpenCLDevice best(OpenCLDevice _deviceLhs, OpenCLDevice _deviceRhs);
+ }
+
+ public static DeviceComparitor Best = new DeviceComparitor(){
+ @Override public OpenCLDevice best(OpenCLDevice _deviceLhs, OpenCLDevice _deviceRhs) {
+ if (_deviceLhs.getType() != _deviceRhs.getType()) {
+ if (_deviceLhs.getType() == TYPE.GPU) {
+ return (_deviceLhs);
+ } else {
+ return (_deviceRhs);
+ }
+ }
+ if (_deviceLhs.maxComputeUnits > _deviceRhs.maxComputeUnits) {
+ return (_deviceLhs);
+ } else {
+ return (_deviceRhs);
+ }
+
+ }
+ };
+
+ public static DeviceFilter FirstGPU = new DeviceFilter(){
+ @Override public boolean match(OpenCLDevice _device) {
+ return (_device.getType() == OpenCLDevice.TYPE.GPU);
+ }
+ };
+
+ public static DeviceFilter FirstCPU = new DeviceFilter(){
+ @Override public boolean match(OpenCLDevice _device) {
+ return (_device.getType() == OpenCLDevice.TYPE.CPU);
+ }
+ };
+
+ public static <T extends OpenCL<T>> T first(Class<T> _interface, DeviceFilter _deviceFilter) {
+ OpenCLDevice device = null;
+ for (OpenCLPlatform p : OpenCLPlatform.getPlatforms()) {
+ for (OpenCLDevice d : p.getDevices()) {
+ if (_deviceFilter.match(d)) {
+ device = d;
+ break;
+ }
+ }
+ if (device != null) {
+ break;
+ }
+ }
+ return (device.create(_interface));
+ }
+
+ public static <T extends OpenCL<T>> T best(Class<T> _interface, DeviceComparitor _deviceComparitor) {
+ OpenCLDevice device = null;
+ for (OpenCLPlatform p : OpenCLPlatform.getPlatforms()) {
+ for (OpenCLDevice d : p.getDevices()) {
+ if (device == null) {
+ device = d;
+ } else {
+ device = _deviceComparitor.best(device, d);
+ }
+ }
+ }
+ return (device.create(_interface));
+ }
+
+ public static <T extends OpenCL<T>> T firstGPU(Class<T> _interface) {
+ return (first(_interface, FirstGPU));
+ }
+
+ public static <T extends OpenCL<T>> T firstCPU(Class<T> _interface) {
+ return (first(_interface, FirstCPU));
+ }
+
+ public static <T extends OpenCL<T>> T best(Class<T> _interface) {
+ return (best(_interface, Best));
+ }
+
+ public OpenCLProgram createProgram(String source) {
+ return (OpenCLJNI.getJNI().createProgram(this, source));
+ }
+
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLJNI.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLJNI.java
new file mode 100644
index 0000000000000000000000000000000000000000..185604f5d9ca736eeeddb0a27e144b69711d15f3
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLJNI.java
@@ -0,0 +1,88 @@
+package com.amd.aparapi;
+
+import java.util.List;
+
+public class OpenCLJNI{
+
+ static {
+ String arch = System.getProperty("os.arch");
+
+ String libName = null;
+ try {
+
+ if (arch.equals("amd64") || arch.equals("x86_64")) {
+ libName = "aparapi_x86_64";
+ } else if (arch.equals("x86") || arch.equals("i386")) {
+ libName = "aparapi_x86";
+ }
+ if (libName != null) {
+ System.out.println("loading " + libName);
+ Runtime.getRuntime().loadLibrary(libName);
+ } else {
+ System.out.println("Expected property os.arch to contain amd64 or x86 but found " + arch
+ + " don't know which library to load.");
+ }
+ } catch (UnsatisfiedLinkError e) {
+ System.out
+ .println("Check your environment. Failed to load aparapi native library "
+ + libName
+ + " or possibly failed to locate opencl native library (opencl.dll/opencl.so). Ensure that both are in your PATH (windows) or in LD_LIBRARY_PATH (linux).");
+
+ }
+
+ Runtime.getRuntime().loadLibrary("aparapi_x86_64");
+ }
+
+ static final OpenCLJNI jni = new OpenCLJNI();
+
+ public static OpenCLJNI getJNI() {
+ return (jni);
+ }
+
+ public final static long INT_BIT = 1 << 0;
+
+ public final static long FLOAT_BIT = 1 << 1;
+
+ public final static long DOUBLE_BIT = 1 << 2;
+
+ public final static long SHORT_BIT = 1 << 3;
+
+ public final static long ARRAY_BIT = 1 << 4;
+
+ public final static long GLOBAL_BIT = 1 << 5;
+
+ public final static long LOCAL_BIT = 1 << 6;
+
+ public final static long CONST_BIT = 1 << 7;
+
+ public final static long PRIMITIVE_BIT = 1 << 8;
+
+ public final static long LONG_BIT = 1 << 9;
+
+ public final static long READONLY_BIT = 1 << 10;
+
+ public final static long WRITEONLY_BIT = 1 << 11;
+
+ public final static long READWRITE_BIT = 1 << 12;
+
+ public final static long MEM_DIRTY_BIT = 1 << 13;
+
+ public final static long MEM_COPY_BIT = 1 << 14;
+
+ public final static long MEM_ENQUEUED_BIT = 1 << 15;
+
+ public final static long ARG_BIT = 1 << 16;
+
+ native public List<OpenCLPlatform> getPlatforms();
+
+ native public OpenCLProgram createProgram(OpenCLDevice context, String openCLSource);
+
+ native public OpenCLKernel createKernel(OpenCLProgram program, String kernelName, List<OpenCLArg> args);
+
+ native public void invoke(OpenCLKernel openCLKernel, Object[] args);
+
+ native public void remap(OpenCLProgram program, OpenCLMem mem, long address);
+
+ native public void getMem(OpenCLProgram program, OpenCLMem mem);
+
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLKernel.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLKernel.java
new file mode 100644
index 0000000000000000000000000000000000000000..74a65a2fb69fa324f4470bb069637b9e5f2bd902
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLKernel.java
@@ -0,0 +1,33 @@
+package com.amd.aparapi;
+
+import java.util.List;
+
+public class OpenCLKernel{
+ private OpenCLArg[] args;
+
+ private long kernelId;
+
+ private OpenCLProgram program;
+
+ private String name;
+
+ OpenCLKernel(long _kernelId, OpenCLProgram _program, String _name, List<OpenCLArg> _args) {
+ kernelId = _kernelId;
+ program = _program;
+ name = _name;
+ args = _args.toArray(new OpenCLArg[0]);
+ for (OpenCLArg arg : args) {
+ arg.kernel = this;
+ }
+ }
+
+ public String getName() {
+ return name;
+ }
+
+ public void invoke(Object[] _args) {
+ OpenCLJNI.getJNI().invoke(this, _args);
+
+ }
+
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLMem.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLMem.java
new file mode 100644
index 0000000000000000000000000000000000000000..2a937f2f69f035902824504e739c962c793bd98f
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLMem.java
@@ -0,0 +1,19 @@
+package com.amd.aparapi;
+
+public class OpenCLMem{
+ long bits;
+
+ int sizeInBytes;
+
+ long readOnlyMemId;
+
+ long writeOnlyMemId;
+
+ long readWriteMemId;
+
+ long address;
+
+ Object instance;
+
+ OpenCLProgram program;
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLPlatform.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLPlatform.java
new file mode 100644
index 0000000000000000000000000000000000000000..1080c3acd96d62821038ecf91b8a3ac0afc85cf1
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLPlatform.java
@@ -0,0 +1,37 @@
+package com.amd.aparapi;
+
+import java.util.ArrayList;
+import java.util.List;
+
+public class OpenCLPlatform{
+ private long platformId;
+
+ private String version;
+
+ private String vendor;
+
+ private List<OpenCLDevice> devices = new ArrayList<OpenCLDevice>();
+
+ OpenCLPlatform(long _platformId, String _version, String _vendor) {
+ platformId = _platformId;
+ version = _version;
+ vendor = _vendor;
+ }
+
+ public String toString() {
+ return ("PlatformId " + platformId + "\nName:" + vendor + "\nVersion:" + version);
+ }
+
+ public void add(OpenCLDevice device) {
+ devices.add(device);
+ }
+
+ public List<OpenCLDevice> getDevices() {
+ return (devices);
+ }
+
+ public static List<OpenCLPlatform> getPlatforms() {
+ return (OpenCLJNI.getJNI().getPlatforms());
+ }
+
+}
diff --git a/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLProgram.java b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLProgram.java
new file mode 100644
index 0000000000000000000000000000000000000000..a42c257747e8431697d2ffb6b8277b28b7df4328
--- /dev/null
+++ b/com.amd.aparapi/src/java/com/amd/aparapi/OpenCLProgram.java
@@ -0,0 +1,64 @@
+package com.amd.aparapi;
+
+import java.util.HashMap;
+import java.util.List;
+import java.util.Map;
+
+public class OpenCLProgram{
+ private long programId;
+
+ private long queueId;
+
+ private long contextId;
+
+ private OpenCLDevice device;
+
+ private String source;
+
+ private String log;
+
+ OpenCLProgram(long _programId, long _queueId, long _contextId, OpenCLDevice _device, String _source, String _log) {
+ programId = _programId;
+ queueId = _queueId;
+ contextId = _contextId;
+ device = _device;
+ source = _source;
+ log = _log;
+ }
+
+ public OpenCLDevice getDevice() {
+ return device;
+ }
+
+ public OpenCLKernel createKernel(String _kernelName, List<OpenCLArg> args) {
+ return (OpenCLJNI.getJNI().createKernel(this, _kernelName, args));
+ }
+
+ private Map<Object, OpenCLMem> instanceToMem = new HashMap<Object, OpenCLMem>();
+
+ private Map<Long, OpenCLMem> addressToMem = new HashMap<Long, OpenCLMem>();
+
+ public synchronized OpenCLMem getMem(Object _instance, long _address) {
+ OpenCLMem mem = instanceToMem.get(_instance);
+ if (mem == null) {
+ mem = addressToMem.get(_instance);
+ if (mem != null) {
+ System.out.println("object has been moved, we need to remap the buffer");
+ OpenCLJNI.getJNI().remap(this, mem, _address);
+ }
+ }
+ return (mem);
+ }
+
+ public synchronized void add(OpenCLMem _mem) {
+
+ instanceToMem.put(_mem.instance, _mem);
+ addressToMem.put(_mem.address, _mem);
+ }
+
+ public synchronized void remaped(OpenCLMem _mem, long _oldAddress) {
+ addressToMem.remove(_oldAddress);
+ addressToMem.put(_mem.address, _mem);
+ }
+
+}
diff --git a/samples/extension/.classpath b/samples/extension/.classpath
new file mode 100644
index 0000000000000000000000000000000000000000..2b3d42947b0a9f028e1acf3ee7d6963e71b1003c
--- /dev/null
+++ b/samples/extension/.classpath
@@ -0,0 +1,8 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<classpath>
+ <classpathentry kind="src" path="src"/>
+ <classpathentry kind="con" path="org.eclipse.jdt.launching.JRE_CONTAINER"/>
+ <classpathentry kind="con" path="org.eclipse.jdt.junit.JUNIT_CONTAINER/4"/>
+ <classpathentry combineaccessrules="false" kind="src" path="/com.amd.aparapi"/>
+ <classpathentry kind="output" path="classes"/>
+</classpath>
diff --git a/samples/extension/.project b/samples/extension/.project
new file mode 100644
index 0000000000000000000000000000000000000000..1e2c59dc69eca99384cd5999f987cdb332a70bdb
--- /dev/null
+++ b/samples/extension/.project
@@ -0,0 +1,17 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<projectDescription>
+ <name>extension</name>
+ <comment></comment>
+ <projects>
+ </projects>
+ <buildSpec>
+ <buildCommand>
+ <name>org.eclipse.jdt.core.javabuilder</name>
+ <arguments>
+ </arguments>
+ </buildCommand>
+ </buildSpec>
+ <natures>
+ <nature>org.eclipse.jdt.core.javanature</nature>
+ </natures>
+</projectDescription>
diff --git a/samples/extension/build.xml b/samples/extension/build.xml
new file mode 100644
index 0000000000000000000000000000000000000000..f6287f1e78aea4bb20054809ad963977e4fd57d3
--- /dev/null
+++ b/samples/extension/build.xml
@@ -0,0 +1,24 @@
+<?xml version="1.0"?>
+
+<project name="extension" default="build" basedir=".">
+ <target name="build" depends="clean">
+ <mkdir dir="classes"/>
+ <javac srcdir="src" destdir="classes" debug="on" includeantruntime="false" >
+ <classpath>
+ <pathelement path="../../com.amd.aparapi/dist/aparapi.jar"/>
+ </classpath>
+ </javac>
+ <copy todir="classes/com\amd\aparapi\sample\extension" file="src\com\amd\aparapi\sample\extension\mandel.cl"/>
+ <copy todir="classes/com\amd\aparapi\sample\extension" file="src\com\amd\aparapi\sample\extension\mandel2.cl"/>
+ <copy todir="classes/com\amd\aparapi\sample\extension" file="src\com\amd\aparapi\sample\extension\squarer.cl"/>
+ <jar jarfile="${ant.project.name}.jar" basedir="classes"/>
+
+ </target>
+
+ <target name="clean">
+ <delete dir="classes"/>
+ <delete file="${ant.project.name}.jar"/>
+ </target>
+
+
+</project>
diff --git a/samples/extension/fft.bat b/samples/extension/fft.bat
new file mode 100644
index 0000000000000000000000000000000000000000..fe6454299113a11360e2fe5756a20a11b22557db
--- /dev/null
+++ b/samples/extension/fft.bat
@@ -0,0 +1,9 @@
+java ^
+ -Djava.library.path=../../com.amd.aparapi.jni/dist ^
+ -Dcom.amd.aparapi.executionMode=%1 ^
+ -Dcom.amd.aparapi.enableProfiling=false ^
+ -Dcom.amd.aparapi.enableShowGeneratedOpenCL=true ^
+ -classpath ../../com.amd.aparapi/dist/aparapi.jar;extension.jar ^
+ com.amd.aparapi.sample.extension.FFTExample
+
+
diff --git a/samples/extension/mandel.bat b/samples/extension/mandel.bat
new file mode 100644
index 0000000000000000000000000000000000000000..7cbefb993303d83a6b245c06ab6037f602f70864
--- /dev/null
+++ b/samples/extension/mandel.bat
@@ -0,0 +1,9 @@
+java ^
+ -Djava.library.path=../../com.amd.aparapi.jni/dist ^
+ -Dcom.amd.aparapi.executionMode=%1 ^
+ -Dcom.amd.aparapi.enableProfiling=false ^
+ -Dcom.amd.aparapi.enableShowGeneratedOpenCL=true ^
+ -classpath ../../com.amd.aparapi/dist/aparapi.jar;extension.jar ^
+ com.amd.aparapi.sample.extension.MandelExample
+
+
diff --git a/samples/extension/square.bat b/samples/extension/square.bat
new file mode 100644
index 0000000000000000000000000000000000000000..01fd7b2a6955648741cb2691617985648a563de0
--- /dev/null
+++ b/samples/extension/square.bat
@@ -0,0 +1,9 @@
+java ^
+ -Djava.library.path=../../com.amd.aparapi.jni/dist ^
+ -Dcom.amd.aparapi.executionMode=%1 ^
+ -Dcom.amd.aparapi.enableProfiling=false ^
+ -Dcom.amd.aparapi.enableShowGeneratedOpenCL=true ^
+ -classpath ../../com.amd.aparapi/dist/aparapi.jar;extension.jar ^
+ com.amd.aparapi.sample.extension.SquareExample
+
+
diff --git a/samples/extension/src/com/amd/aparapi/sample/extension/MandelExample.java b/samples/extension/src/com/amd/aparapi/sample/extension/MandelExample.java
new file mode 100644
index 0000000000000000000000000000000000000000..e788939bdaa2a231ba82e725ad14e6d69f0f9413
--- /dev/null
+++ b/samples/extension/src/com/amd/aparapi/sample/extension/MandelExample.java
@@ -0,0 +1,495 @@
+/*
+Copyright (c) 2010-2011, Advanced Micro Devices, Inc.
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without modification, are permitted provided that the
+following conditions are met:
+
+Redistributions of source code must retain the above copyright notice, this list of conditions and the following
+disclaimer.
+
+Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following
+disclaimer in the documentation and/or other materials provided with the distribution.
+
+Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products
+derived from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES,
+INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export
+laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730 through
+774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of the EAR,
+you hereby certify that, except pursuant to a license granted by the United States Department of Commerce Bureau of
+Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export Administration
+Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in Country Groups D:1,
+E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2) export to Country Groups
+D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced direct product is subject
+to national security controls as identified on the Commerce Control List (currently found in Supplement 1 to Part 774
+of EAR). For the most current Country Group listings, or for additional information about the EAR or your obligations
+under those regulations, please refer to the U.S. Bureau of Industry and Security's website at http://www.bis.doc.gov/.
+
+*/
+
+package com.amd.aparapi.sample.extension;
+
+import java.awt.BorderLayout;
+import java.awt.Dimension;
+import java.awt.FlowLayout;
+import java.awt.Graphics;
+import java.awt.Point;
+import java.awt.event.ItemEvent;
+import java.awt.event.ItemListener;
+import java.awt.event.MouseAdapter;
+import java.awt.event.MouseEvent;
+import java.awt.event.WindowAdapter;
+import java.awt.event.WindowEvent;
+import java.awt.image.BufferedImage;
+import java.awt.image.DataBufferInt;
+import java.util.concurrent.BrokenBarrierException;
+import java.util.concurrent.CyclicBarrier;
+
+import javax.swing.JComboBox;
+import javax.swing.JComponent;
+import javax.swing.JFrame;
+import javax.swing.JLabel;
+import javax.swing.JPanel;
+import javax.swing.JTextField;
+
+import com.amd.aparapi.OpenCL;
+import com.amd.aparapi.OpenCLAdapter;
+import com.amd.aparapi.OpenCLDevice;
+import com.amd.aparapi.Range;
+
+/**
+ * An example Aparapi application which displays a view of the Mandelbrot set and lets the user zoom in to a particular point.
+ *
+ * When the user clicks on the view, this example application will zoom in to the clicked point and zoom out there after.
+ * On GPU, additional computing units will offer a better viewing experience. On the other hand on CPU, this example
+ * application might suffer with sub-optimal frame refresh rate as compared to GPU.
+ *
+ * @author gfrost
+ *
+ */
+
+@OpenCL.Resource("com/amd/aparapi/sample/extension/mandel2.cl") interface MandelBrot extends OpenCL<MandelBrot>{
+ MandelBrot createMandleBrot(//
+ Range range,//
+ @Arg("scale") float scale, //
+ @Arg("offsetx") float offsetx, //
+ @Arg("offsety") float offsety, //
+ @GlobalWriteOnly("rgb") int[] rgb
+
+ );
+}
+
+class JavaMandelBrot extends OpenCLAdapter<MandelBrot> implements MandelBrot{
+ final int MAX_ITERATIONS = 64;
+
+ final int pallette[] = new int[] {
+ -65536,
+ -59392,
+ -53248,
+ -112640,
+ -106752,
+ -166144,
+ -160256,
+ -219904,
+ -279552,
+ -339200,
+ -399104,
+ -985344,
+ -2624000,
+ -4197376,
+ -5770496,
+ -7343872,
+ -8851712,
+ -10425088,
+ -11932928,
+ -13375232,
+ -14817792,
+ -16260096,
+ -16719602,
+ -16720349,
+ -16721097,
+ -16721846,
+ -16722595,
+ -16723345,
+ -16724351,
+ -16725102,
+ -16726110,
+ -16727119,
+ -16728129,
+ -16733509,
+ -16738889,
+ -16744269,
+ -16749138,
+ -16754006,
+ -16758619,
+ -16762976,
+ -16767077,
+ -16771178,
+ -16774767,
+ -16514932,
+ -15662970,
+ -14942079,
+ -14221189,
+ -13631371,
+ -13107088,
+ -12648342,
+ -12320669,
+ -11992995,
+ -11796393,
+ -11665328,
+ -11993019,
+ -12386248,
+ -12845011,
+ -13303773,
+ -13762534,
+ -14286830,
+ -14745588,
+ -15269881,
+ -15728637,
+ -16252927,
+ 0
+ };
+
+ @Override public MandelBrot createMandleBrot(Range range, float scale, float offsetx, float offsety, int[] rgb) {
+
+ int width = range.getGlobalSize(0);
+ int height = range.getGlobalSize(1);
+ for (int gridy = 0; gridy < height; gridy++) {
+ for (int gridx = 0; gridx < width; gridx++) {
+ float x = ((((float) (gridx) * scale) - ((scale / 2.0f) * (float) width)) / (float) width) + offsetx;
+ float y = ((((float) (gridy) * scale) - ((scale / 2.0f) * (float) height)) / (float) height) + offsety;
+ int count = 0;
+ float zx = x;
+ float zy = y;
+ float new_zx = 0.0f;
+ for (; count < MAX_ITERATIONS && ((zx * zx) + (zy * zy)) < 8.0f; count++) {
+ new_zx = ((zx * zx) - (zy * zy)) + x;
+ zy = ((2.0f * zx) * zy) + y;
+ zx = new_zx;
+ }
+ rgb[gridx + gridy * width] = pallette[count];
+
+ }
+ }
+ return (this);
+ }
+
+}
+
+class JavaMandelBrotMultiThread extends OpenCLAdapter<MandelBrot> implements MandelBrot{
+ final int MAX_ITERATIONS = 64;
+
+ final int pallette[] = new int[] {
+ -65536,
+ -59392,
+ -53248,
+ -112640,
+ -106752,
+ -166144,
+ -160256,
+ -219904,
+ -279552,
+ -339200,
+ -399104,
+ -985344,
+ -2624000,
+ -4197376,
+ -5770496,
+ -7343872,
+ -8851712,
+ -10425088,
+ -11932928,
+ -13375232,
+ -14817792,
+ -16260096,
+ -16719602,
+ -16720349,
+ -16721097,
+ -16721846,
+ -16722595,
+ -16723345,
+ -16724351,
+ -16725102,
+ -16726110,
+ -16727119,
+ -16728129,
+ -16733509,
+ -16738889,
+ -16744269,
+ -16749138,
+ -16754006,
+ -16758619,
+ -16762976,
+ -16767077,
+ -16771178,
+ -16774767,
+ -16514932,
+ -15662970,
+ -14942079,
+ -14221189,
+ -13631371,
+ -13107088,
+ -12648342,
+ -12320669,
+ -11992995,
+ -11796393,
+ -11665328,
+ -11993019,
+ -12386248,
+ -12845011,
+ -13303773,
+ -13762534,
+ -14286830,
+ -14745588,
+ -15269881,
+ -15728637,
+ -16252927,
+ 0
+ };
+
+ @Override public MandelBrot createMandleBrot(final Range range, final float scale, final float offsetx, final float offsety,
+ final int[] rgb) {
+
+ final int width = range.getGlobalSize(0);
+ final int height = range.getGlobalSize(1);
+ final int threadCount = 8;
+ Thread[] threads = new Thread[threadCount];
+ final CyclicBarrier barrier = new CyclicBarrier(threadCount + 1);
+ for (int thread = 0; thread < threadCount; thread++) {
+ final int threadId = thread;
+ final int groupHeight = height / threadCount;
+ (threads[threadId] = new Thread(new Runnable(){
+ public void run() {
+ for (int gridy = threadId * groupHeight; gridy < (threadId + 1) * groupHeight; gridy++) {
+ for (int gridx = 0; gridx < width; gridx++) {
+ float x = ((((float) (gridx) * scale) - ((scale / 2.0f) * (float) width)) / (float) width) + offsetx;
+ float y = ((((float) (gridy) * scale) - ((scale / 2.0f) * (float) height)) / (float) height) + offsety;
+ int count = 0;
+ float zx = x;
+ float zy = y;
+ float new_zx = 0.0f;
+ for (; count < MAX_ITERATIONS && ((zx * zx) + (zy * zy)) < 8.0f; count++) {
+ new_zx = ((zx * zx) - (zy * zy)) + x;
+ zy = ((2.0f * zx) * zy) + y;
+ zx = new_zx;
+ }
+ rgb[gridx + gridy * width] = pallette[count];
+ }
+ }
+ try {
+ barrier.await();
+ } catch (InterruptedException e) {
+ // TODO Auto-generated catch block
+ e.printStackTrace();
+ } catch (BrokenBarrierException e) {
+ // TODO Auto-generated catch block
+ e.printStackTrace();
+ }
+ }
+ })).start();
+ }
+ try {
+ barrier.await();
+ } catch (InterruptedException e) {
+ // TODO Auto-generated catch block
+ e.printStackTrace();
+ } catch (BrokenBarrierException e) {
+ // TODO Auto-generated catch block
+ e.printStackTrace();
+ }
+ return (this);
+ }
+
+}
+
+public class MandelExample{
+
+ /** User selected zoom-in point on the Mandelbrot view. */
+ public static volatile Point to = null;
+
+ public static MandelBrot mandelBrot = null;
+
+ public static MandelBrot gpuMandelBrot = null;
+
+ public static MandelBrot javaMandelBrot = null;
+
+ public static MandelBrot javaMandelBrotMultiThread = null;
+
+ // new JavaMandelBrot();
+ //new JavaMandelBrotMultiThread();
+ @SuppressWarnings("serial") public static void main(String[] _args) {
+
+ JFrame frame = new JFrame("MandelBrot");
+
+ /** Width of Mandelbrot view. */
+ final int width = 768;
+
+ /** Height of Mandelbrot view. */
+ final int height = 768;
+
+ /** Mandelbrot image height. */
+ final Range range = Range.create2D(width, height);
+
+ /** Image for Mandelbrot view. */
+ final BufferedImage image = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB);
+
+ final Object framePaintedDoorBell = new Object();
+ JComponent viewer = new JComponent(){
+ @Override public void paintComponent(Graphics g) {
+
+ g.drawImage(image, 0, 0, width, height, this);
+ synchronized (framePaintedDoorBell) {
+ framePaintedDoorBell.notify();
+ }
+ }
+ };
+
+ // Set the size of JComponent which displays Mandelbrot image
+ viewer.setPreferredSize(new Dimension(width, height));
+
+ final Object userClickDoorBell = new Object();
+
+ // Mouse listener which reads the user clicked zoom-in point on the Mandelbrot view
+ viewer.addMouseListener(new MouseAdapter(){
+ @Override public void mouseClicked(MouseEvent e) {
+ to = e.getPoint();
+ synchronized (userClickDoorBell) {
+ userClickDoorBell.notify();
+ }
+ }
+ });
+
+ JPanel controlPanel = new JPanel(new FlowLayout());
+
+ final String[] choices = new String[] {
+ "Java Sequential",
+ "Java Threads",
+ "GPU OpenCL"
+ };
+
+ final JComboBox startButton = new JComboBox(choices);
+
+ startButton.addItemListener(new ItemListener(){
+ @Override public void itemStateChanged(ItemEvent e) {
+ String item = (String) startButton.getSelectedItem();
+
+ if (item.equals(choices[2])) {
+ mandelBrot = gpuMandelBrot;
+ } else if (item.equals(choices[0])) {
+ mandelBrot = javaMandelBrot;
+ } else if (item.equals(choices[1])) {
+ mandelBrot = javaMandelBrotMultiThread;
+ }
+ }
+
+ });
+ controlPanel.add(startButton);
+
+ controlPanel.add(new JLabel("FPS"));
+ final JTextField framesPerSecondTextField = new JTextField("0", 5);
+
+ controlPanel.add(framesPerSecondTextField);
+
+ // Swing housework to create the frame
+ frame.getContentPane().add(viewer);
+ frame.getContentPane().add(controlPanel, BorderLayout.SOUTH);
+ frame.pack();
+ frame.setLocationRelativeTo(null);
+ frame.setVisible(true);
+
+ // Extract the underlying RGB buffer from the image.
+ // Pass this to the kernel so it operates directly on the RGB buffer of the image
+
+ final int[] imageRgb = ((DataBufferInt) image.getRaster().getDataBuffer()).getData();
+
+ /** Mutable values of scale, offsetx and offsety so that we can modify the zoom level and position of a view. */
+ float scale = .0f;
+
+ float offsetx = .0f;
+
+ float offsety = .0f;
+ gpuMandelBrot = OpenCLDevice.best(MandelBrot.class);
+ javaMandelBrot = new JavaMandelBrot();
+ javaMandelBrotMultiThread = new JavaMandelBrotMultiThread();
+ mandelBrot = javaMandelBrot;
+ float defaultScale = 3f;
+ scale = defaultScale;
+ offsetx = -1f;
+ offsety = 0f;
+
+ mandelBrot.createMandleBrot(range, scale, offsetx, offsety, imageRgb);
+ viewer.repaint();
+
+ // Window listener to dispose Kernel resources on user exit.
+ frame.addWindowListener(new WindowAdapter(){
+ public void windowClosing(WindowEvent _windowEvent) {
+ // mandelBrot.dispose();
+ System.exit(0);
+ }
+ });
+
+ while (true) {
+ // Wait for the user to click somewhere
+ while (to == null) {
+ synchronized (userClickDoorBell) {
+ try {
+ userClickDoorBell.wait();
+ } catch (InterruptedException ie) {
+ ie.getStackTrace();
+ }
+ }
+ }
+
+ float x = -1f;
+ float y = 0f;
+ float tox = (float) (to.x - width / 2) / width * scale;
+ float toy = (float) (to.y - height / 2) / height * scale;
+
+ // This is how many frames we will display as we zoom in and out.
+ int frames = 128;
+ long startMillis = System.currentTimeMillis();
+ int frameCount = 0;
+ for (int sign = -1; sign < 2; sign += 2) {
+ for (int i = 0; i < frames - 4; i++) {
+ scale = scale + sign * defaultScale / frames;
+ x = x - sign * (tox / frames);
+ y = y - sign * (toy / frames);
+ offsetx = x;
+ offsety = y;
+ mandelBrot.createMandleBrot(range, scale, offsetx, offsety, imageRgb);
+ viewer.repaint();
+ synchronized (framePaintedDoorBell) {
+ try {
+ framePaintedDoorBell.wait();
+ } catch (InterruptedException ie) {
+ ie.getStackTrace();
+ }
+ }
+ frameCount++;
+ long endMillis = System.currentTimeMillis();
+ long elapsedMillis = endMillis - startMillis;
+ if (elapsedMillis > 1000) {
+ framesPerSecondTextField.setText("" + frameCount * 1000 / elapsedMillis);
+ frameCount = 0;
+ startMillis = endMillis;
+ }
+ }
+ }
+
+ // Reset zoom-in point.
+ to = null;
+
+ }
+
+ }
+
+}
diff --git a/samples/extension/src/com/amd/aparapi/sample/extension/SquareExample.java b/samples/extension/src/com/amd/aparapi/sample/extension/SquareExample.java
new file mode 100644
index 0000000000000000000000000000000000000000..33708495813cab9edb8708e444cc4c8fa09e9f8e
--- /dev/null
+++ b/samples/extension/src/com/amd/aparapi/sample/extension/SquareExample.java
@@ -0,0 +1,68 @@
+package com.amd.aparapi.sample.extension;
+
+import com.amd.aparapi.OpenCL;
+import com.amd.aparapi.OpenCLDevice;
+import com.amd.aparapi.Range;
+
+public class SquareExample{
+
+ interface Squarer extends OpenCL<Squarer>{
+ @Kernel("{\n"//
+ + " const size_t id = get_global_id(0);\n"//
+ + " out[id] = in[id]*in[id];\n"//
+ + "}\n")//
+ public Squarer square(//
+ Range _range,//
+ @GlobalReadWrite("in") float[] in,//
+ @GlobalReadWrite("out") float[] out);
+ }
+
+ @OpenCL.Resource("com/amd/aparapi/sample/extension/squarer.cl") interface SquarerWithResource extends
+ OpenCL<SquarerWithResource>{
+
+ public SquarerWithResource square(//
+ Range _range,//
+ @GlobalReadWrite("in") float[] in,//
+ @GlobalReadWrite("out") float[] out);
+ }
+
+ @OpenCL.Source("\n"//
+ + "__kernel void square (\n" //
+ + " __global float *in,\n"//
+ + " __global float *out\n" + "){\n"//
+ + " const size_t id = get_global_id(0);\n"//
+ + " out[id] = in[id]*in[id];\n"//
+ + "}\n") interface SquarerWithSource extends OpenCL<SquarerWithSource>{
+
+ public SquarerWithSource square(//
+ Range _range,//
+ @GlobalReadOnly("in") float[] in,//
+ @GlobalWriteOnly("out") float[] out);
+ }
+
+ public static void main(String[] args) {
+
+ int size = 32;
+ int[] v = new int[size];
+ float[] in = new float[size];
+ for (int i = 0; i < size; i++) {
+ in[i] = i;
+ }
+ float[] out = new float[size];
+ Range range = Range.create(size);
+
+ SquarerWithResource squarer = OpenCLDevice.firstGPU(SquarerWithResource.class);
+ squarer.square(range, in, out);
+
+ for (int i = 0; i < size; i++) {
+ System.out.println(in[i] + " " + out[i]);
+ }
+
+ squarer.square(range, out, in);
+
+ for (int i = 0; i < size; i++) {
+ System.out.println(in[i] + " " + out[i]);
+ }
+ }
+
+}
diff --git a/samples/extension/src/com/amd/aparapi/sample/extension/mandel.cl b/samples/extension/src/com/amd/aparapi/sample/extension/mandel.cl
new file mode 100644
index 0000000000000000000000000000000000000000..b6c1da0137393bb537fc39148f9ab59abdf64328
--- /dev/null
+++ b/samples/extension/src/com/amd/aparapi/sample/extension/mandel.cl
@@ -0,0 +1,91 @@
+#define MAX_ITERATIONS 64
+
+__constant const int pallette[]={
+ -65536,
+ -59392,
+ -53248,
+ -112640,
+ -106752,
+ -166144,
+ -160256,
+ -219904,
+ -279552,
+ -339200,
+ -399104,
+ -985344,
+ -2624000,
+ -4197376,
+ -5770496,
+ -7343872,
+ -8851712,
+ -10425088,
+ -11932928,
+ -13375232,
+ -14817792,
+ -16260096,
+ -16719602,
+ -16720349,
+ -16721097,
+ -16721846,
+ -16722595,
+ -16723345,
+ -16724351,
+ -16725102,
+ -16726110,
+ -16727119,
+ -16728129,
+ -16733509,
+ -16738889,
+ -16744269,
+ -16749138,
+ -16754006,
+ -16758619,
+ -16762976,
+ -16767077,
+ -16771178,
+ -16774767,
+ -16514932,
+ -15662970,
+ -14942079,
+ -14221189,
+ -13631371,
+ -13107088,
+ -12648342,
+ -12320669,
+ -11992995,
+ -11796393,
+ -11665328,
+ -11993019,
+ -12386248,
+ -12845011,
+ -13303773,
+ -13762534,
+ -14286830,
+ -14745588,
+ -15269881,
+ -15728637,
+ -16252927,
+ 0
+};
+
+__kernel void createMandleBrot(
+ float scale,
+ float offsetx,
+ float offsety,
+ __global int *rgb
+){
+ int gid = get_global_id(0) + get_global_id(1)*get_global_size(0);
+ float x = ((((float)(get_global_id(0)) * scale) - ((scale / 2.0f) * (float)get_global_size(0))) / (float)get_global_size(0)) + offsetx;
+ float y = ((((float)(get_global_id(1)) * scale) - ((scale / 2.0f) * (float)get_global_size(1))) / (float)get_global_size(1)) + offsety;
+ int count = 0;
+ float zx = x;
+ float zy = y;
+ float new_zx = 0.0f;
+ for (; count<MAX_ITERATIONS && ((zx * zx) + (zy * zy))<8.0f; count++){
+ new_zx = ((zx * zx) - (zy * zy)) + x;
+ zy = ((2.0f * zx) * zy) + y;
+ zx = new_zx;
+ }
+ rgb[gid] = pallette[count];
+}
+
diff --git a/samples/extension/src/com/amd/aparapi/sample/extension/mandel2.cl b/samples/extension/src/com/amd/aparapi/sample/extension/mandel2.cl
new file mode 100644
index 0000000000000000000000000000000000000000..e7043427233a395c51eaff67ff6e5fbeb05eafd7
--- /dev/null
+++ b/samples/extension/src/com/amd/aparapi/sample/extension/mandel2.cl
@@ -0,0 +1,95 @@
+#define MAX_ITERATIONS 64
+
+__constant const int pallette[]={
+ -65536,
+ -59392,
+ -53248,
+ -112640,
+ -106752,
+ -166144,
+ -160256,
+ -219904,
+ -279552,
+ -339200,
+ -399104,
+ -985344,
+ -2624000,
+ -4197376,
+ -5770496,
+ -7343872,
+ -8851712,
+ -10425088,
+ -11932928,
+ -13375232,
+ -14817792,
+ -16260096,
+ -16719602,
+ -16720349,
+ -16721097,
+ -16721846,
+ -16722595,
+ -16723345,
+ -16724351,
+ -16725102,
+ -16726110,
+ -16727119,
+ -16728129,
+ -16733509,
+ -16738889,
+ -16744269,
+ -16749138,
+ -16754006,
+ -16758619,
+ -16762976,
+ -16767077,
+ -16771178,
+ -16774767,
+ -16514932,
+ -15662970,
+ -14942079,
+ -14221189,
+ -13631371,
+ -13107088,
+ -12648342,
+ -12320669,
+ -11992995,
+ -11796393,
+ -11665328,
+ -11993019,
+ -12386248,
+ -12845011,
+ -13303773,
+ -13762534,
+ -14286830,
+ -14745588,
+ -15269881,
+ -15728637,
+ -16252927,
+ 0
+};
+
+#define WIDTH get_global_size(0)
+#define HEIGHT get_global_size(1)
+#define X get_global_id(0)
+#define Y get_global_id(1)
+
+__kernel void createMandleBrot(
+ float scale,
+ float offsetx,
+ float offsety,
+ __global int *rgb
+ ){
+ float x = ((((float)(X) * scale) - ((scale / 2.0f) * (float)WIDTH)) / (float)WIDTH) + offsetx;
+ float y = ((((float)(Y) * scale) - ((scale / 2.0f) * (float)HEIGHT)) / (float)HEIGHT) + offsety;
+ float zx = x;
+ float zy = y;
+ float new_zx = 0.0f;
+ int count = 0;
+ for (; count<MAX_ITERATIONS && ((zx * zx) + (zy * zy))<8.0f; count++){
+ new_zx = ((zx * zx) - (zy * zy)) + x;
+ zy = ((2.0f * zx) * zy) + y;
+ zx = new_zx;
+ }
+ rgb[X + Y*WIDTH] = pallette[count];
+}
+
diff --git a/samples/extension/src/com/amd/aparapi/sample/extension/squarer.cl b/samples/extension/src/com/amd/aparapi/sample/extension/squarer.cl
new file mode 100644
index 0000000000000000000000000000000000000000..7169cfc1e09db0d099cf6f80de2f7592bc277aec
--- /dev/null
+++ b/samples/extension/src/com/amd/aparapi/sample/extension/squarer.cl
@@ -0,0 +1,5 @@
+__kernel void square( __global float *in, __global float *out){
+ const size_t id = get_global_id(0);
+ out[id] = in[id]*in[id];
+}
+