private boolean test_float_3D(){
  CUarray array=new CUarray();
  CUDA_ARRAY3D_DESCRIPTOR ad=new CUDA_ARRAY3D_DESCRIPTOR();
  ad.Format=CU_AD_FORMAT_FLOAT;
  ad.Width=sizeX;
  ad.Height=sizeY;
  ad.Depth=sizeZ;
  ad.NumChannels=1;
  cuArray3DCreate(array,ad);
  CUDA_MEMCPY3D copy=new CUDA_MEMCPY3D();
  copy.srcMemoryType=CUmemorytype.CU_MEMORYTYPE_HOST;
  copy.srcHost=Pointer.to(input_float_3D);
  copy.srcPitch=sizeX * Sizeof.FLOAT;
  copy.srcHeight=sizeY;
  copy.dstMemoryType=CUmemorytype.CU_MEMORYTYPE_ARRAY;
  copy.dstArray=array;
  copy.dstHeight=sizeX;
  copy.WidthInBytes=sizeX * Sizeof.FLOAT;
  copy.Height=sizeY;
  copy.Depth=sizeZ;
  cuMemcpy3D(copy);
  CUtexref texref=new CUtexref();
  cuModuleGetTexRef(texref,module,"texture_float_3D");
  cuTexRefSetFilterMode(texref,CU_TR_FILTER_MODE_LINEAR);
  cuTexRefSetAddressMode(texref,0,CU_TR_ADDRESS_MODE_CLAMP);
  cuTexRefSetAddressMode(texref,1,CU_TR_ADDRESS_MODE_CLAMP);
  cuTexRefSetAddressMode(texref,2,CU_TR_ADDRESS_MODE_CLAMP);
  cuTexRefSetFlags(texref,CU_TRSF_NORMALIZED_COORDINATES);
  cuTexRefSetFormat(texref,CU_AD_FORMAT_FLOAT,1);
  cuTexRefSetArray(texref,array,CU_TRSA_OVERRIDE_FORMAT);
  CUdeviceptr dOutput=new CUdeviceptr();
  cuMemAlloc(dOutput,Sizeof.FLOAT * 1);
  CUfunction function=new CUfunction();
  cuModuleGetFunction(function,module,"test_float_3D");
  Pointer kernelParameters=Pointer.to(Pointer.to(dOutput),Pointer.to(new float[]{posX}),Pointer.to(new float[]{posY}),Pointer.to(new float[]{posZ}));
  cuLaunchKernel(function,1,1,1,1,1,1,0,null,kernelParameters,null);
  cuCtxSynchronize();
  float hOutput[]=new float[1];
  cuMemcpyDtoH(Pointer.to(hOutput),dOutput,Sizeof.FLOAT * 1);
  log("Result float  3D " + Arrays.toString(hOutput));
  float expected[]=new float[]{3.5f};
  boolean passed=Arrays.equals(hOutput,expected);
  log("Test   float  3D " + (passed ? "PASSED" : "FAILED"));
  cuArrayDestroy(array);
  cuMemFree(dOutput);
  return passed;
}

private boolean test_float_3D ( ) { CUarray array = new CUarray ( ) ; CUDA_ARRAY3D_DESCRIPTOR ad = new CUDA_ARRAY3D_DESCRIPTOR ( ) ; ad . Format = CU_AD_FORMAT_FLOAT ; ad . Width = sizeX ; ad . Height = sizeY ; ad . Depth = sizeZ ; ad . NumChannels = 1 ; cuArray3DCreate ( array , ad ) ; CUDA_MEMCPY3D copy = new CUDA_MEMCPY3D ( ) ; copy . srcMemoryType = CUmemorytype . CU_MEMORYTYPE_HOST ; copy . srcHost = Pointer . to ( input_float_3D ) ; copy . srcPitch = sizeX * Sizeof . FLOAT ; copy . srcHeight = sizeY ; copy . dstMemoryType = CUmemorytype . CU_MEMORYTYPE_ARRAY ; copy . dstArray = array ; copy . dstHeight = sizeX ; copy . WidthInBytes = sizeX * Sizeof . FLOAT ; copy . Height = sizeY ; copy . Depth = sizeZ ; cuMemcpy3D ( copy ) ; CUtexref texref = new CUtexref ( ) ; cuModuleGetTexRef ( texref , module , "texture_float_3D" ) ; cuTexRefSetFilterMode ( texref , CU_TR_FILTER_MODE_LINEAR ) ; cuTexRefSetAddressMode ( texref , 0 , CU_TR_ADDRESS_MODE_CLAMP ) ; cuTexRefSetAddressMode ( texref , 1 , CU_TR_ADDRESS_MODE_CLAMP ) ; cuTexRefSetAddressMode ( texref , 2 , CU_TR_ADDRESS_MODE_CLAMP ) ; cuTexRefSetFlags ( texref , CU_TRSF_NORMALIZED_COORDINATES ) ; cuTexRefSetFormat ( texref , CU_AD_FORMAT_FLOAT , 1 ) ; cuTexRefSetArray ( texref , array , CU_TRSA_OVERRIDE_FORMAT ) ; CUdeviceptr dOutput = new CUdeviceptr ( ) ; cuMemAlloc ( dOutput , Sizeof . FLOAT * 1 ) ; CUfunction function = new CUfunction ( ) ; cuModuleGetFunction ( function , module , "test_float_3D" ) ; Pointer kernelParameters = Pointer . to ( Pointer . to ( dOutput ) , Pointer . to ( new float [ ] { posX } ) , Pointer . to ( new float [ ] { posY } ) , Pointer . to ( new float [ ] { posZ } ) ) ; cuLaunchKernel ( function , 1 , 1 , 1 , 1 , 1 , 1 , 0 , null , kernelParameters , null ) ; cuCtxSynchronize ( ) ; float hOutput [ ] = new float [ 1 ] ; cuMemcpyDtoH ( Pointer . to ( hOutput ) , dOutput , Sizeof . FLOAT * 1 ) ; log ( "Result float  3D " + Arrays . toString ( hOutput ) ) ; float expected [ ] = new float [ ] { 3.5f } ; boolean passed = Arrays . equals ( hOutput , expected ) ; log ( "Test   float  3D " + ( passed ? "PASSED" : "FAILED" ) ) ; cuArrayDestroy ( array ) ; cuMemFree ( dOutput ) ; return passed ; }
test the 3d float texture access
