cnugteren / cltune Goto Github PK

[----------] Testing kernel gemm_fast
device compiler error/warning: (277:0) : error : require any typed expression except arrays, structures containing arrays, sampler types, and structures containing sampler types
(293:0) : error : require any typed expression except arrays, structures containing arrays, sampler types, and structures containing sampler types
(321:0) : error : require a matching typed expression
(388:0) : error : require a struct/union or vector typed expression
(389:0) : error : require a struct/union or vector typed expression
(494:0) : error : function: 'LocalToPrivateA' hasn't the corresponding declaration
(502:0) : error : function: 'LocalToPrivateB' hasn't the corresponding declaration
(509:0) : error : function: 'MultiplyAccumulate' hasn't the corresponding declaration
(520:0) : error : function: 'StoreResults' hasn't the corresponding declaration

[   FAILED ] Kernel gemm_fast failed
[   FAILED ]   catched exception: device compiler error/warning occurred ^^

[   FAILED ] gemm_fast;      0.0 ms;   MWG 32;  NWG 128;   KWG 32; MDIMC 16; NDIMC 16; MDIMA 16; NDIMB 32;    KWI 2;    VWM 2;    VWN 2;   STRM 0;   STRN 0;     SA 1;     SB 1;PRECISION 32;

[==========] Initializing on platform 0 device 0
[==========] Device name: 'Intel(R) Xeon(R) CPU E5-1620 v2 @ 3.70GHz' (OpenCL 1.2 AMD-APP (1912.5))

[----------] Testing reference gemm_reference
[ RUN      ] Running gemm_reference
[       OK ] Completed gemm_reference (10915.9 ms) - 1 out of 1

[----------] Testing kernel gemm_fast
[ RUN      ] Running gemm_fast
[   FAILED ] Kernel gemm_fast failed
[   FAILED ]   catched exception: Internal OpenCL error: -54
[  WARNING ] Results differ: L2 norm is 1.90e+08
[   FAILED ] gemm_fast;      0.0 ms;  MWG 128;  NWG 128;   KWG 32; MDIMC 32;  NDIMC 8; MDIMA 32; NDIMB 16;    KWI 8;    VWM 1;    VWN 1;   STRM 1;   STRN 0;     SA 1;     SB 1;PRECISION 32;
[ RUN      ] Running gemm_fast
[   FAILED ] Kernel gemm_fast failed
[   FAILED ]   catched exception: Internal OpenCL error: -54
[  WARNING ] Results differ: L2 norm is 1.90e+08
[   FAILED ] gemm_fast;      0.0 ms;  MWG 128;   NWG 64;   KWG 32; MDIMC 32; NDIMC 16; MDIMA 32; NDIMB 16;    KWI 8;    VWM 1;    VWN 4;   STRM 1;   STRN 0;     SA 1;     SB 1;PRECISION 32;
[ RUN      ] Running gemm_fast
[   FAILED ] Kernel gemm_fast failed
[   FAILED ]   catched exception: Internal OpenCL error: -54
[  WARNING ] Results differ: L2 norm is 1.90e+08
[   FAILED ] gemm_fast;      0.0 ms;  MWG 128;   NWG 64;   KWG 32; MDIMC 32; NDIMC 16; MDIMA 32; NDIMB 32;    KWI 8;    VWM 1;    VWN 2;   STRM 1;   STRN 0;     SA 0;     SB 1;PRECISION 32;

...
[database] Saving database to './scripts/database/database.json'
[database] Calculating the best results per device/kernel...
[database] Calculating the default values...
[ERROR] Entries for a single kernel with multiple argument values: ARM;GPU;32;xgemm;Xgemm

__kernel void saxpy( const          size_t   N,
                     const          float    a,
                     const __global float*   x,
                           __global float*   y
                   )
{
  for( int w = 0; w < WPT; ++w ) {
    const int id = w * get_global_size(0)
                     + get_global_id(0);

    y[ id ] += a * x[ id ];
  }
}

// CLTune program for tuning saxpy
int main()
{
  // Set kernel and input size
  const auto   saxpy = /* path to kernel */
  const size_t N       = /* input size     */

  // Declare input data
  float a;
  auto  vec_x = std::vector<float>( N );
  auto  vec_y = std::vector<float>( N );

  // Initialize random number generator
  const auto random_seed = std::chrono::system_clock::now().time_since_epoch().count();
  std::default_random_engine generator(static_cast<unsigned int>(random_seed));
  std::uniform_real_distribution<float> distribution(-2.0f, 2.0f);

  // Fill input data
  a = distribution(generator);
  for (auto &item: vec_x) { item = distribution(generator); }
  for (auto &item: vec_y) { item = distribution(generator); }

  // Initializes the tuner (platform 1, device 0)
  cltune::Tuner tuner( 1, 0 );
  tuner.UseAnnealing( 1.0f / 2048.0f, 4.0 );
  auto id = tuner.AddKernel( saxpy, "saxpy", { N }, { 1 } );
  
  // Set tuning-parameters
  tuner.AddParameter(id, "LS", {1, ... , N} );
  tuner.AddParameter(id, "WPT", {1, ... ,N} );

  // Sets constraints on parameters
  auto DividesN       = []( std::vector<size_t> v ){ return  N % v[0]         == 0; };
  auto DividesNDivWPT = []( std::vector<size_t> v ){ return (N / v[0]) % v[1] == 0; };
 
  tuner.AddConstraint(id, DividesN      , {"WPT"}        );
  tuner.AddConstraint(id, DividesNDivWPT, {"WPT", "LS"} );

  // Modifies the thread-sizes (both global and local) based on the parameters
  tuner.MulLocalSize(id, {"LS"} );
  tuner.DivGlobalSize(id, {"WPT" } );

  // Sets the function's arguments.
  tuner.AddArgumentScalar( N );
  tuner.AddArgumentScalar( a );
  tuner.AddArgumentInput(vec_x);
  tuner.AddArgumentOutput(vec_y);

  // Starts the tuner
  tuner.Tune();
  tuner.PrintToFile("output.csv");
}

Device compiler error/warning: <unknown>:0:0: in function conv void (i32, i32, float addrspace(1)*, float addrspace(2)*, float addrspace(1)*): unsupported call to function LoadLocalPlusHalo
[   FAILED ] Kernel conv failed

./sample_conv
[==========] Initializing on platform 0 device 0
[==========] Device name: 'AMD TONGA (DRM 3.8.0 / 4.9.11-1-ARCH, LLVM 3.9.1)' (OpenCL 1.1 Mesa 17.0.0)
[----------] Testing reference conv_reference
[ RUN      ] Running conv_reference
[       OK ] Completed conv_reference (65.7 ms) - 1 out of 1
[----------] Testing kernel conv
[ RUN      ] Running conv
[       OK ] Completed conv (86.1 ms) - 1 out of 22
[ RUN      ] Running conv
[       OK ] Completed conv (91.9 ms) - 2 out of 22
[ RUN      ] Running conv
[       OK ] Completed conv (67.3 ms) - 3 out of 22