private boolean test_float4_2D(){
  CUarray array=new CUarray();
  CUDA_ARRAY_DESCRIPTOR ad=new CUDA_ARRAY_DESCRIPTOR();
  ad.Format=CU_AD_FORMAT_FLOAT;
  ad.Width=sizeX;
  ad.Height=sizeY;
  ad.NumChannels=4;
  cuArrayCreate(array,ad);
  CUDA_MEMCPY2D copyHD=new CUDA_MEMCPY2D();
  copyHD.srcMemoryType=CUmemorytype.CU_MEMORYTYPE_HOST;
  copyHD.srcHost=Pointer.to(input_float4_2D);
  copyHD.srcPitch=sizeX * Sizeof.FLOAT * 4;
  copyHD.dstMemoryType=CUmemorytype.CU_MEMORYTYPE_ARRAY;
  copyHD.dstArray=array;
  copyHD.WidthInBytes=sizeX * Sizeof.FLOAT * 4;
  copyHD.Height=sizeY;
  cuMemcpy2D(copyHD);
  CUtexref texref=new CUtexref();
  cuModuleGetTexRef(texref,module,"texture_float4_2D");
  cuTexRefSetFilterMode(texref,CU_TR_FILTER_MODE_LINEAR);
  cuTexRefSetAddressMode(texref,0,CU_TR_ADDRESS_MODE_CLAMP);
  cuTexRefSetAddressMode(texref,1,CU_TR_ADDRESS_MODE_CLAMP);
  cuTexRefSetFlags(texref,CU_TRSF_NORMALIZED_COORDINATES);
  cuTexRefSetFormat(texref,CU_AD_FORMAT_FLOAT,4);
  cuTexRefSetArray(texref,array,CU_TRSA_OVERRIDE_FORMAT);
  CUdeviceptr dOutput=new CUdeviceptr();
  cuMemAlloc(dOutput,Sizeof.FLOAT * 4);
  CUfunction function=new CUfunction();
  cuModuleGetFunction(function,module,"test_float4_2D");
  Pointer kernelParameters=Pointer.to(Pointer.to(dOutput),Pointer.to(new float[]{posX}),Pointer.to(new float[]{posY}));
  cuLaunchKernel(function,1,1,1,1,1,1,0,null,kernelParameters,null);
  cuCtxSynchronize();
  float hOutput[]=new float[4];
  cuMemcpyDtoH(Pointer.to(hOutput),dOutput,Sizeof.FLOAT * 4);
  log("Result float4 2D " + Arrays.toString(hOutput));
  float expected[]=new float[]{1.5f,1.5f,1.5f,1.5f};
  boolean passed=Arrays.equals(hOutput,expected);
  log("Test   float4 2D " + (passed ? "PASSED" : "FAILED"));
  cuArrayDestroy(array);
  cuMemFree(dOutput);
  return passed;
}

private boolean test_float4_2D ( ) { CUarray array = new CUarray ( ) ; CUDA_ARRAY_DESCRIPTOR ad = new CUDA_ARRAY_DESCRIPTOR ( ) ; ad . Format = CU_AD_FORMAT_FLOAT ; ad . Width = sizeX ; ad . Height = sizeY ; ad . NumChannels = 4 ; cuArrayCreate ( array , ad ) ; CUDA_MEMCPY2D copyHD = new CUDA_MEMCPY2D ( ) ; copyHD . srcMemoryType = CUmemorytype . CU_MEMORYTYPE_HOST ; copyHD . srcHost = Pointer . to ( input_float4_2D ) ; copyHD . srcPitch = sizeX * Sizeof . FLOAT * 4 ; copyHD . dstMemoryType = CUmemorytype . CU_MEMORYTYPE_ARRAY ; copyHD . dstArray = array ; copyHD . WidthInBytes = sizeX * Sizeof . FLOAT * 4 ; copyHD . Height = sizeY ; cuMemcpy2D ( copyHD ) ; CUtexref texref = new CUtexref ( ) ; cuModuleGetTexRef ( texref , module , "texture_float4_2D" ) ; cuTexRefSetFilterMode ( texref , CU_TR_FILTER_MODE_LINEAR ) ; cuTexRefSetAddressMode ( texref , 0 , CU_TR_ADDRESS_MODE_CLAMP ) ; cuTexRefSetAddressMode ( texref , 1 , CU_TR_ADDRESS_MODE_CLAMP ) ; cuTexRefSetFlags ( texref , CU_TRSF_NORMALIZED_COORDINATES ) ; cuTexRefSetFormat ( texref , CU_AD_FORMAT_FLOAT , 4 ) ; cuTexRefSetArray ( texref , array , CU_TRSA_OVERRIDE_FORMAT ) ; CUdeviceptr dOutput = new CUdeviceptr ( ) ; cuMemAlloc ( dOutput , Sizeof . FLOAT * 4 ) ; CUfunction function = new CUfunction ( ) ; cuModuleGetFunction ( function , module , "test_float4_2D" ) ; Pointer kernelParameters = Pointer . to ( Pointer . to ( dOutput ) , Pointer . to ( new float [ ] { posX } ) , Pointer . to ( new float [ ] { posY } ) ) ; cuLaunchKernel ( function , 1 , 1 , 1 , 1 , 1 , 1 , 0 , null , kernelParameters , null ) ; cuCtxSynchronize ( ) ; float hOutput [ ] = new float [ 4 ] ; cuMemcpyDtoH ( Pointer . to ( hOutput ) , dOutput , Sizeof . FLOAT * 4 ) ; log ( "Result float4 2D " + Arrays . toString ( hOutput ) ) ; float expected [ ] = new float [ ] { 1.5f , 1.5f , 1.5f , 1.5f } ; boolean passed = Arrays . equals ( hOutput , expected ) ; log ( "Test   float4 2D " + ( passed ? "PASSED" : "FAILED" ) ) ; cuArrayDestroy ( array ) ; cuMemFree ( dOutput ) ; return passed ; }
test the 2d float4 texture access
