prepared for separate compilation, for now merged (3d5ddc28) · Commits · Elphel / imagej-elphel

src/main/java/com/elphel/imagej/gpu/GPUTileProcessor.java

+154 −66

Original line number	Diff line number	Diff line
		@@ -31,6 +31,7 @@ import static jcuda.driver.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABI
		// Uses code by Marco Hutter - http://www.jcuda.org
		import static jcuda.driver.CUjitInputType.CU_JIT_INPUT_LIBRARY;
		import static jcuda.driver.CUjitInputType.CU_JIT_INPUT_PTX;
		import static jcuda.driver.CUjit_option.CU_JIT_LOG_VERBOSE;
		import static jcuda.driver.JCudaDriver.cuCtxCreate;
		import static jcuda.driver.JCudaDriver.cuCtxSynchronize;
		import static jcuda.driver.JCudaDriver.cuDeviceGet;
		@@ -62,6 +63,7 @@ import java.io.IOException;
		import java.nio.charset.StandardCharsets;
		import java.nio.file.Files;
		import java.nio.file.Paths;
		import java.util.Random;
		import java.util.concurrent.atomic.AtomicInteger;

		import com.elphel.imagej.tileprocessor.DttRad2;
		@@ -89,6 +91,10 @@ public class GPUTileProcessor {
		String LIBRARY_PATH = "/usr/local/cuda/targets/x86_64-linux/lib/libcudadevrt.a"; // linux
		static String GPU_RESOURCE_DIR = "kernels";
		static String [] GPU_KERNEL_FILES = {"dtt8x8.cuh","TileProcessor.cuh"};
		// "*" - generated defines, first index - separately compiled unit
		// static String [][] GPU_SRC_FILES = {{"","dtt8x8.h","dtt8x8.cu"},{"","dtt8x8.h","TileProcessor.cuh"}};
		static String [][] GPU_SRC_FILES = {{"*","dtt8x8.h","dtt8x8.cu","TileProcessor.cuh"}};
		// static String [][] GPU_SRC_FILES = {{"*","dtt8x8.cuh","TileProcessor.cuh"}};
		static String GPU_CONVERT_CORRECT_TILES_NAME = "convert_correct_tiles"; // name in C code
		static String GPU_IMCLT_RBG_NAME = "imclt_rbg"; // name in C code
		static String GPU_CORRELATE2D_NAME = "correlate2D"; // name in C code
		@@ -295,6 +301,41 @@ public class GPUTileProcessor {
		return new PointerWithAddress(p).getAddress();
		}

		private String getTpDefines() {
		return"#define JCUDA\n"+
		"#define DTT_SIZE_LOG2 " + DTT_SIZE_LOG2+"\n"+
		"#define THREADSX " + THREADSX+"\n"+
		"#define NUM_CAMS " + NUM_CAMS+"\n"+
		"#define NUM_PAIRS " + NUM_PAIRS+"\n"+
		"#define NUM_COLORS " + NUM_COLORS+"\n"+
		"#define IMG_WIDTH " + IMG_WIDTH+"\n"+
		"#define IMG_HEIGHT " + IMG_HEIGHT+"\n"+
		"#define KERNELS_HOR " + KERNELS_HOR+"\n"+
		"#define KERNELS_VERT " + KERNELS_VERT+"\n"+
		"#define KERNELS_LSTEP " + KERNELS_LSTEP+"\n"+
		"#define THREADS_PER_TILE " + THREADS_PER_TILE+"\n"+
		"#define TILES_PER_BLOCK " + TILES_PER_BLOCK+"\n"+
		"#define CORR_THREADS_PER_TILE " + CORR_THREADS_PER_TILE+"\n"+
		"#define CORR_TILES_PER_BLOCK " + CORR_TILES_PER_BLOCK+"\n"+
		"#define TEXTURE_THREADS_PER_TILE " + TEXTURE_THREADS_PER_TILE+"\n"+
		"#define TEXTURE_TILES_PER_BLOCK " + TEXTURE_TILES_PER_BLOCK+"\n"+
		"#define IMCLT_THREADS_PER_TILE " + IMCLT_THREADS_PER_TILE+"\n"+
		"#define IMCLT_TILES_PER_BLOCK " + IMCLT_TILES_PER_BLOCK+"\n"+
		"#define CORR_NTILE_SHIFT " + CORR_NTILE_SHIFT+"\n"+
		"#define CORR_PAIRS_MASK " + CORR_PAIRS_MASK+"\n"+
		"#define CORR_TEXTURE_BIT " + CORR_TEXTURE_BIT+"\n"+
		"#define TASK_CORR_BITS " + TASK_CORR_BITS+"\n"+
		"#define TASK_TEXTURE_N_BIT " + TASK_TEXTURE_N_BIT+"\n"+
		"#define TASK_TEXTURE_E_BIT " + TASK_TEXTURE_E_BIT+"\n"+
		"#define TASK_TEXTURE_S_BIT " + TASK_TEXTURE_S_BIT+"\n"+
		"#define TASK_TEXTURE_W_BIT " + TASK_TEXTURE_W_BIT+"\n"+
		"#define LIST_TEXTURE_BIT " + LIST_TEXTURE_BIT+"\n"+
		"#define CORR_OUT_RAD " + CORR_OUT_RAD+"\n" +
		"#define FAT_ZERO_WEIGHT " + FAT_ZERO_WEIGHT+"\n"+
		"#define THREADS_DYNAMIC_BITS " + THREADS_DYNAMIC_BITS+"\n";

		}

		public GPUTileProcessor(String cuda_project_directory) throws IOException
		{

		@@ -326,40 +367,38 @@ public class GPUTileProcessor {
		// When using just Eclipse resources - it does not notice that the file
		// was edited (happens frequently during kernel development).
		ClassLoader classLoader = getClass().getClassLoader();
		String kernelSource =
		"#define JCUDA\n"+
		"#define DTT_SIZE_LOG2 " + DTT_SIZE_LOG2+"\n"+
		"#define THREADSX " + THREADSX+"\n"+
		"#define NUM_CAMS " + NUM_CAMS+"\n"+
		"#define NUM_PAIRS " + NUM_PAIRS+"\n"+
		"#define NUM_COLORS " + NUM_COLORS+"\n"+
		"#define IMG_WIDTH " + IMG_WIDTH+"\n"+
		"#define IMG_HEIGHT " + IMG_HEIGHT+"\n"+
		"#define KERNELS_HOR " + KERNELS_HOR+"\n"+
		"#define KERNELS_VERT " + KERNELS_VERT+"\n"+
		"#define KERNELS_LSTEP " + KERNELS_LSTEP+"\n"+
		"#define THREADS_PER_TILE " + THREADS_PER_TILE+"\n"+
		"#define TILES_PER_BLOCK " + TILES_PER_BLOCK+"\n"+
		"#define CORR_THREADS_PER_TILE " + CORR_THREADS_PER_TILE+"\n"+
		"#define CORR_TILES_PER_BLOCK " + CORR_TILES_PER_BLOCK+"\n"+
		"#define TEXTURE_THREADS_PER_TILE " + TEXTURE_THREADS_PER_TILE+"\n"+
		"#define TEXTURE_TILES_PER_BLOCK " + TEXTURE_TILES_PER_BLOCK+"\n"+
		"#define IMCLT_THREADS_PER_TILE " + IMCLT_THREADS_PER_TILE+"\n"+
		"#define IMCLT_TILES_PER_BLOCK " + IMCLT_TILES_PER_BLOCK+"\n"+
		"#define CORR_NTILE_SHIFT " + CORR_NTILE_SHIFT+"\n"+
		"#define CORR_PAIRS_MASK " + CORR_PAIRS_MASK+"\n"+
		"#define CORR_TEXTURE_BIT " + CORR_TEXTURE_BIT+"\n"+
		"#define TASK_CORR_BITS " + TASK_CORR_BITS+"\n"+
		"#define TASK_TEXTURE_N_BIT " + TASK_TEXTURE_N_BIT+"\n"+
		"#define TASK_TEXTURE_E_BIT " + TASK_TEXTURE_E_BIT+"\n"+
		"#define TASK_TEXTURE_S_BIT " + TASK_TEXTURE_S_BIT+"\n"+
		"#define TASK_TEXTURE_W_BIT " + TASK_TEXTURE_W_BIT+"\n"+
		"#define LIST_TEXTURE_BIT " + LIST_TEXTURE_BIT+"\n"+
		"#define CORR_OUT_RAD " + CORR_OUT_RAD+"\n" +
		"#define FAT_ZERO_WEIGHT " + FAT_ZERO_WEIGHT+"\n"+
		"#define THREADS_DYNAMIC_BITS " + THREADS_DYNAMIC_BITS+"\n";

		String [] kernelSources = new String[GPU_SRC_FILES.length];

		for (int cunit = 0; cunit < kernelSources.length; cunit++) {
		kernelSources[cunit] = ""; // use StringBuffer?
		for (String src_file:GPU_SRC_FILES[cunit]) {
		if (src_file.contentEquals("*")) {
		kernelSources[cunit] += getTpDefines();
		}else {
		File file = null;
		if ((cuda_project_directory == null) \|\| cuda_project_directory.isEmpty()) {
		file = new File(classLoader.getResource(GPU_RESOURCE_DIR+"/"+src_file).getFile());
		System.out.println("Loading resource "+file);
		} else {
		File src_dir = new File(cuda_project_directory, "src");
		file = new File(src_dir.getPath(), src_file);
		System.out.println("Loading resource "+file);
		}
		System.out.println(file.getAbsolutePath());
		String cuFileName = file.getAbsolutePath(); // /home/eyesis/workspace-python3/nvidia_dct8x8/src/dtt8x8.cuh";// "dtt8x8.cuh";
		String sourceFile = readFileAsString(cuFileName); // readResourceAsString(cuFileName);
		if (sourceFile == null) {
		String msg = "Could not read the kernel source code from "+cuFileName;
		IJ.showMessage("Error", msg);
		new IllegalArgumentException (msg);
		}
		kernelSources[cunit] += sourceFile;
		}
		}
		}
		/*
		String kernelSource = getTpDefines();
		for (String src_file:GPU_KERNEL_FILES) {
		File file = null;
		if ((cuda_project_directory == null) \|\| cuda_project_directory.isEmpty()) {
		@@ -379,8 +418,9 @@ public class GPUTileProcessor {
		new IllegalArgumentException (msg);
		}
		kernelSource += sourceFile;

		}
		*/

		// Create the kernel functions (first - just test)
		String [] func_names = {
		GPU_CONVERT_CORRECT_TILES_NAME,
		@@ -388,7 +428,7 @@ public class GPUTileProcessor {
		GPU_CORRELATE2D_NAME,
		GPU_TEXTURES_NAME,
		GPU_RBGA_NAME};
		CUfunction[] functions = createFunctions(kernelSource,
		CUfunction[] functions = createFunctions(kernelSources,
		func_names,
		capability); // on my - 75

		@@ -711,6 +751,7 @@ public class GPUTileProcessor {
		double xc = woi.x + rx - 0.5;
		double yc = woi.y + ry - 0.5;
		boolean dbg1 = false; // true;
		double dbg_frac = 0.0; // 0.25;
		boolean [] mask = new boolean[tilesX*tilesY];
		int num_tiles = 0;
		for (int ty = woi.y; ty < (woi.y +woi.height); ty++) {
		@@ -725,6 +766,35 @@ public class GPUTileProcessor {
		}
		}
		}
		if (dbg_frac > 0) {
		Random rnd = new Random(0);
		int num_final = (int) Math.round(num_tiles * (1.0 - dbg_frac));
		while (num_tiles > num_final) {
		int tx = woi.x + rnd.nextInt(woi.width);
		int ty = woi.y + rnd.nextInt(woi.height);
		int indx = ty * tilesX + tx;
		if (mask[indx]) {
		mask[indx] = false;
		num_tiles--;
		}
		}
		// filter out with no neighbors
		for (int indx = 0; indx < mask.length; indx++) if (mask[indx]) {
		int ix = indx % tilesX;
		int iy = indx / tilesX;
		int num_neib = 0;
		if ((ix > 0) && mask[indx-1]) num_neib++;
		if ((ix < (tilesX-1)) && mask[indx+1]) num_neib++;
		if ((iy > 0) && mask[indx-tilesX]) num_neib++;
		if ((iy < (tilesY-1)) && mask[indx+tilesX]) num_neib++;
		if (num_neib == 0) {
		mask[indx] = false;
		num_tiles--;
		}
		}
		//nextInt(int bound)
		}

		if (dbg1) {
		// mask[(woi.y-1) * tilesX + (woi.x-1)] = true;
		mask[(woi.y+woi.height) * tilesX + (woi.x+woi.width)] = true;
		@@ -1332,13 +1402,17 @@ public class GPUTileProcessor {

		// private static CUfunction [] createFunctions(
		private CUfunction [] createFunctions(
		String sourceCode,
		String [] sourceCodeUnits,
		String [] kernelNames,
		int capability
		) throws IOException
		{
		CUfunction [] functions = new CUfunction [kernelNames.length];
		byte[][] ptxDataUnits = new byte [sourceCodeUnits.length][];
		boolean OK = false;
		// for (String sourceCode: sourceCodeUnits) {
		for (int cunit = 0; cunit < ptxDataUnits.length; cunit++) {
		String sourceCode = sourceCodeUnits[cunit];
		// Use the NVRTC to create a program by compiling the source code
		nvrtcProgram program = new nvrtcProgram();
		nvrtcCreateProgram( program, sourceCode, null, 0, null, null);
		@@ -1366,19 +1440,33 @@ public class GPUTileProcessor {
		String[] ptx = new String[1];
		nvrtcGetPTX(program, ptx);
		nvrtcDestroyProgram(program);
		byte[] ptxData = ptx[0].getBytes();
		// byte[] ptxData = ptx[0].getBytes();
		ptxDataUnits[cunit] = ptx[0].getBytes();
		System.out.println("ptxDataUnits["+cunit+"].length="+ptxDataUnits[cunit].length);
		// System.out.println( ptx[0]);
		}
		JITOptions jitOptions = new JITOptions();
		jitOptions.putInt(CU_JIT_LOG_VERBOSE, 1);
		CUlinkState state = new CUlinkState();
		cuLinkCreate(jitOptions, state);
		cuLinkAddFile(state, CU_JIT_INPUT_LIBRARY, LIBRARY_PATH, jitOptions);

		System.out.println("ptxData.length="+ptxData.length);
		// System.out.println( ptx[0]);
		for (int cunit = 0; cunit < ptxDataUnits.length; cunit++) {
		// cuLinkAddData(state, CU_JIT_INPUT_PTX, Pointer.to(ptxData), ptxData.length, "input.ptx", jitOptions); // CUDA_ERROR_INVALID_PTX
		cuLinkAddData(state, CU_JIT_INPUT_PTX, Pointer.to(ptxDataUnits[cunit]), ptxDataUnits[cunit].length, "input"+cunit+".ptx", jitOptions); // CUDA_ERROR_INVALID_PTX
		// cuLinkAddData(state, CU_JIT_INPUT_PTX, Pointer.to(ptxDataUnits[cunit]), ptxDataUnits[cunit].length, "input.ptx", jitOptions); // CUDA_ERROR_INVALID_PTX
		}
		// cuLinkAddFile(state, CU_JIT_INPUT_LIBRARY, LIBRARY_PATH, jitOptions);

		cuLinkAddData(state, CU_JIT_INPUT_PTX, Pointer.to(ptxData), ptxData.length, "input.ptx", jitOptions); // CUDA_ERROR_INVALID_PTX
		long size[] = { 0 };
		Pointer image = new Pointer();
		cuLinkComplete(state, image, size);
		JCudaDriver.setExceptionsEnabled(false);
		int cuda_result = cuLinkComplete(state, image, size);
		System.out.println("cuLinkComplete() -> "+cuda_result);

		JCudaDriver.setExceptionsEnabled(true);


		module = new CUmodule();
		cuModuleLoadDataEx(module, image, 0, new int[0], Pointer.to(new int[0]));
		cuLinkDestroy(state);

src/main/resources/kernels/TileProcessor.cuh

+166 −943

File changed.

Preview size limit exceeded, changes collapsed.

src/main/resources/kernels/dtt8x8.cuh→src/main/resources/kernels/dtt8x8.cu

+0 −0

File changed and moved.

Preview size limit exceeded, changes collapsed.

src/main/resources/kernels/dtt8x8.h

0 → 100644

+111 −0

Original line number	Diff line number	Diff line
		/**
		**
		** dtt8x8.h
		**
		** Copyright (C) 2018 Elphel, Inc.
		**
		-----------------------------------------------------------------------------
		**
		** dtt8x8.cuh is free software: you can redistribute it and/or modify
		** it under the terms of the GNU General Public License as published by
		** the Free Software Foundation, either version 3 of the License, or
		** (at your option) any later version.
		**
		** This program is distributed in the hope that it will be useful,
		** but WITHOUT ANY WARRANTY; without even the implied warranty of
		** MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
		** GNU General Public License for more details.
		**
		** You should have received a copy of the GNU General Public License
		** along with this program. If not, see <http://www.gnu.org/licenses/>.
		**
		** Additional permission under GNU GPL version 3 section 7
		**
		** If you modify this Program, or any covered work, by linking or
		** combining it with NVIDIA Corporation's CUDA libraries from the
		** NVIDIA CUDA Toolkit (or a modified version of those libraries),
		** containing parts covered by the terms of NVIDIA CUDA Toolkit
		** EULA, the licensors of this Program grant you additional
		** permission to convey the resulting work.
		-----------------------------------------------------------------------------
		*/

		/**
		**************************************************************************
		* \file dtt8x8.h
		* \brief DCT-II, DST-II, DCT-IV and DST-IV for Complex Lapped Transform of 16x16 (stride 8)
		* in GPU
		* This file contains building blocks for the 16x16 stride 8 COmplex Lapped Transform (CLT)
		* implementation. DTT-IV are used for forward and inverse 2D CLT, DTT-II - to convert correlation
		* results from the frequency to pixel domain. DTT-III (inverse of DTT-II) is not implemented
		* here it is used to convert convolution kernels and LPF to the frequency domain - done in
		* software.
		*
		* This file is cpompatible with both runtime and driver API, runtime is used for development
		* with Nvidia Nsight, driver API when calling these kernels from Java
		*/
		#ifndef JCUDA
		#define DTT_SIZE_LOG2 3
		#endif

		#pragma once
		#define DTT_SIZE (1 << DTT_SIZE_LOG2)
		#define DTT_SIZE1 (DTT_SIZE + 1)
		#define DTT_SIZE2 (2 * DTT_SIZE)
		#define DTT_SIZE21 (DTT_SIZE2 + 1)
		#define DTT_SIZE4 (4 * DTT_SIZE)
		#define DTT_SIZE2M1 (DTT_SIZE2 - 1)
		#define BAYER_RED 0
		#define BAYER_BLUE 1
		#define BAYER_GREEN 2
		// assuming GR/BG as now
		#define BAYER_RED_ROW 0
		#define BAYER_RED_COL 1

		#define DTTTEST_BLOCK_WIDTH 32
		#define DTTTEST_BLOCK_HEIGHT 16
		#define DTTTEST_BLK_STRIDE (DTTTEST_BLOCK_WIDTH+1)

		//extern __constant__ float idct_signs[4][4][4];
		//extern __constant__ int imclt_indx9[16];
		//extern __constant__ float HWINDOW2[];

		inline __device__ void dttii_shared_mem_nonortho(float * x0, int inc, int dst_not_dct); // does not scale by y[0] (y[7]) by 1/sqrt[0]
		inline __device__ void dttii_shared_mem(float * x0, int inc, int dst_not_dct); // used in GPU_DTT24_DRV
		inline __device__ void dttiv_shared_mem(float * x0, int inc, int dst_not_dct); // used in GPU_DTT24_DRV
		inline __device__ void dttiv_nodiverg (float * x, int inc, int dst_not_dct); // not used
		inline __device__ void dctiv_nodiverg (float * x0, int inc); // used in TP
		inline __device__ void dstiv_nodiverg (float * x0, int inc); // used in TP

		inline __device__ void dct_ii8 ( float x[8], float y[8]); // x,y point to 8-element arrays each // not used
		inline __device__ void dct_iv8 ( float x[8], float y[8]); // x,y point to 8-element arrays each // not used
		inline __device__ void dst_iv8 ( float x[8], float y[8]); // x,y point to 8-element arrays each // not used
		inline __device__ void _dctii_nrecurs8 ( float x[8], float y[8]); // x,y point to 8-element arrays each // not used
		inline __device__ void _dctiv_nrecurs8 ( float x[8], float y[8]); // x,y point to 8-element arrays each // not used

		// kernels (not used so far)
		#ifdef BBBB
		extern "C" __global__ void GPU_DTT24_DRV(float dst, float src, int src_stride, int dtt_mode);
		#endif// #ifdef BBBB

		//=========================== 2D functions ===============
		extern __device__ void corrUnfoldTile(
		int corr_radius,
		float* qdata0, // [4][DTT_SIZE][DTT_SIZE1], // 4 quadrants of the clt data, rows extended to optimize shared ports
		float* rslt); // [DTT_SIZE2M1][DTT_SIZE2M1]) // 15x15

		extern __device__ void dttii_2d(
		float * clt_corr); // shared memory, [4][DTT_SIZE1][DTT_SIZE]

		extern __device__ void dttiv_color_2d(
		float * clt_tile,
		int color);
		extern __device__ void imclt(
		float * clt_tile, // [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports [4][8][9]
		float * mclt_tile );

		extern __device__ void imclt8threads(
		int do_acc, // 1 - add to previous value, 0 - overwrite
		float * clt_tile, // [4][DTT_SIZE][DTT_SIZE1], // +1 to alternate column ports [4][8][9]
		float * mclt_tile, // [2* DTT_SIZE][DTT_SIZE1+ DTT_SIZE], // +1 to alternate column ports[16][17]
		int debug);

src/main/resources/kernels/test_tp.cu

0 → 100644

+1097 −0

File added.

Preview size limit exceeded, changes collapsed.

Admin message