text:
/* HIT_START
 * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -std=c++11
 * TEST: %t
 * HIT_END
 */

#include <cstdint>
#include "hip/hip_runtime.h"
#include "test_common.h"

// Memory alignment is broken
// Update: with latest changes the aligment is working fine, hence enabled
#define ENABLE_ALIGNMENT_TEST_SMALL_BAR 1

// Packed member atribute broken
#define ENABLE_PACKED_TEST 0

// Update: with latest changes struct class object
// from device is working fine, hence enabled
#define ENABLE_CLASS_OBJ_ACCESS 1

// accessing dynamic/heap memory from device is broken
#define ENABLE_HEAP_MEMORY_ACCESS 0

// Update: with latest changes it's working hence enabled
#define ENABLE_USER_STL 1

// Update: with latest changes it's working hence enabled
#define ENABLE_OUT_OF_ORDER_INITIALIZATION 1

// Direct initialization of struct broken,
// ip_d9 is a pointer, uint_t*, hipLaunchKernelStruct_h9 = {'c', ip_d9};
#define ENABLE_DECLARE_INITIALIZATION_POINTER 0

// Bit fields are broken
#define ENABLE_BIT_FIELDS 0

static const int  BLOCK_DIM_SIZE = 512;

// allocate memory on device and host for result validation
static bool *result_d, *result_h;
static hipError_t hipMallocError = hipMalloc((void**)&result_d,
                                              BLOCK_DIM_SIZE*sizeof(bool));
static hipError_t hipHostMallocError = hipHostMalloc((void**)&result_h,
                                              BLOCK_DIM_SIZE*sizeof(bool));
static hipError_t hipMemsetError = hipMemset(result_d,
                                              false, BLOCK_DIM_SIZE);

static void ResultValidation() {
    hipMemcpy(result_h, result_d, BLOCK_DIM_SIZE*sizeof(bool),
              hipMemcpyDeviceToHost);

    for (int k = 0; k < BLOCK_DIM_SIZE; ++k) {
      HIPASSERT(result_h[k] == true);
    }
    return;
}

// Segregating the reset part as it was causing a problem when i put inside
// ResultValidation() function, the memory was not reset correctly for the
// tests which were disabled.
static void ResetValidationMem() {
    // reset the memory to false to reuse it.
    hipMemset(result_d, false, BLOCK_DIM_SIZE);
    hipMemset(result_h, false, BLOCK_DIM_SIZE);
    return;
}

// This test is to verify Struct with variables
// support, read from device.
typedef struct hipLaunchKernelStruct1 {
  int li;  // local int
  float lf;  // local float
  bool result;  // local bool
} hipLaunchKernelStruct_t1;

// This test is to verify struct with padding, read from device
typedef struct hipLaunchKernelStruct2 {
  char c1;
  long l1;
  char c2;
  long l2;
  bool result;
} hipLaunchKernelStruct_t2;

// This test is to verify struct with padding, read from device
typedef struct hipLaunchKernelStruct3 {
  char bf1;
  char bf2;
  long l1;
  char bf3;
  bool result;
} hipLaunchKernelStruct_t3;

// This test is to verify empty struct
typedef struct hipLaunchKernelStruct4 {
  // empty struct, size will be verified from device side,size 1Byte
} hipLaunchKernelStruct_t4;

// This test is to verify struct with pointer member variable.
typedef struct hipLaunchKernelStruct5 {
  char c1;
  char* cp;  // char pointer
} hipLaunchKernelStruct_t5;

// This test is to verify struct with aligned(8),
// right now it's broken on hcc & hip-clang
typedef struct hipLaunchKernelStruct6 {
  char c1;
  short int si;
} __attribute__((aligned(8)))  hipLaunchKernelStruct_t6;

// This test is to verify struct with aligned(16),
// right now it's brokenon hcc & hip-clang
typedef struct hipLaunchKernelStruct7 {
  char c1;
  short int si;
} __attribute__((aligned(16)))  hipLaunchKernelStruct_t7;

// This test is to verify struct with packed & aligned,
// size should be 4Bytes right now it's broken on hcc & hip-clang
typedef struct hipLaunchKernelStruct8 {
  char c1;
  short int si;
  bool b;
}__attribute__((packed, aligned(4)))  hipLaunchKernelStruct_t8;

// This test is to verify struct with packed, no alignment as Sam suggested
// size should be 4Bytes, right now it's broken on hcc & hip-clang
typedef struct hipLaunchKernelStruct8A {
  char c1;
  short int si;
  bool b;
}__attribute__((packed))  hipLaunchKernelStruct_t8A;

// This test is to verify struct with alignment, no packing as Sam suggested
// size should be 8Bytes as no packing, right now it's broken on hcc & hip-clang
typedef struct hipLaunchKernelStruct8B {
  char c1;
  short int si;
  bool b;
}__attribute__((aligned(8)))  hipLaunchKernelStruct_t8B;

// This test is to verify const struct object
typedef struct hipLaunchKernelStruct9 {
  char c1;
  uint32_t* ip;  // uint pointer
} hipLaunchKernelStruct_t9;

// This test is to verify struct with stdint types, uintN_t
typedef struct hipLaunchKernelStruct10 {
  uint64_t u64;
  uint32_t u32;
  uint8_t u8;
} hipLaunchKernelStruct_t10;

// This test is to verify struct with volatile member
typedef struct hipLaunchKernelStruct11 {
  int i1;
  volatile unsigned int vint;
} hipLaunchKernelStruct_t11;

// This test is to verify struct with simple class object
class base {
 public:
    int i = 0;
    base() {}
};
typedef struct hipLaunchKernelStruct12 {
  base b;
  char c1;
} hipLaunchKernelStruct_t12;

// This test is to verify struct with __device__ func() attribute
typedef struct hipLaunchKernelStruct13 {
  int i1;
  __device__ int getvalue() { return i1; }
} hipLaunchKernelStruct_t13;

// This test is to verify struct with array variable,
// write to from device
typedef struct hipLaunchKernelStruct14 {
  int readint;
  int writeint[BLOCK_DIM_SIZE];  // will write to this from device
} hipLaunchKernelStruct_t14;

// This test is to verify struct with dynamic memory, new int
// the heap memory will be accessed from device
typedef struct hipLaunchKernelStruct15 {
  char c1;
  int* heapmem;  // allocated using hipMalloc()
} hipLaunchKernelStruct_t15;

// This test is to verify simple template struct
template<typename T>
struct hipLaunchKernelStruct_t16 {
  T t1;
};

// This test is to verify simple explicity template struct
template<typename T> struct hipLaunchKernelStruct_t17 {};
template<>  // explicit template
struct hipLaunchKernelStruct_t17<int> {
  int t1;
};

// This test is to verity write to struct memory using __device__ func()
typedef struct hipLaunchKernelStruct18 {
  char c1;
  __device__ void setChar(char c) { c1 = c; }
  __device__ int getChar() { return c1; }
} hipLaunchKernelStruct_t18;

// This test is to verity user defined STL, simple stack implementation
typedef struct stackNode {
    int data;
    stackNode* nextNode = NULL;
} stackNode_t;
typedef struct hipLaunchKernelStruct19 {
  stackNode_t* stack = NULL;
  unsigned int size_ = 0;
  void pushMe(int value) {  // not a device function, setting from host
    stackNode_t* newNode;
    hipMalloc((void**)&newNode, sizeof(stackNode_t));
    hipMemset(&newNode->data, value, sizeof(stackNode_t));
    //newNode->data = value;
    ++size_;
    if (stack == NULL) {
      stack = newNode;
      return;
    }
    stackNode_t* currentHead = stack;
    stack = newNode;
    stack->nextNode = currentHead;
    return;
  }
  __device__ void popMe() {
    stackNode_t* currentHead = stack;
    stack = stack->nextNode;
    --size_;
    // delete currentHead;  // no idea why delete not working
    return;
  }
  int stackSize() {
    return size_;
  }
} hipLaunchKernelStruct_t19;

// This test is to verify out of order initalizer of struct elements
// and access in-order, from device.
typedef struct hipLaunchKernelStruct20 {
  char name;
  int age;
  int rank;
} hipLaunchKernelStruct_t20;

// This test is to verify bit fields operations
// the size should be 1Bytes
typedef struct hipLaunchKernelStruct21 {
  int i : 3;  // limiting bits to 3
  int j : 2;  // limiting bits to 2
} hipLaunchKernelStruct_t21;

// Passing struct to a hipLaunchKernelGGL(),
// read and write into the same struct
__global__ void hipLaunchKernelStructFunc1(
                    hipLaunchKernelStruct_t1 hipLaunchKernelStruct_,
                    bool* result_d1) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d1[x] =  ((hipLaunchKernelStruct_.li == 1)
                      && (hipLaunchKernelStruct_.lf == 1.0)
                      && (hipLaunchKernelStruct_.result == false));
}

// Passing struct to a hipLaunchKernelGGL(), checks padding,
// read and write into the same struct
__global__ void hipLaunchKernelStructFunc2(
                    hipLaunchKernelStruct_t2 hipLaunchKernelStruct_,
                    bool* result_d2) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d2[x] =  ((hipLaunchKernelStruct_.c1 == 'a')
                      && (hipLaunchKernelStruct_.l1 == 1.0)
                      && (hipLaunchKernelStruct_.c2 == 'b')
                      && (hipLaunchKernelStruct_.l2 == 2.0) );
}

// Passing struct to a hipLaunchKernelGGL(), checks padding,
// read and write into the same struct
__global__ void hipLaunchKernelStructFunc3(
                    hipLaunchKernelStruct_t3 hipLaunchKernelStruct_,
                    bool* result_d3) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d3[x] =  ((hipLaunchKernelStruct_.bf1 == 1)
                     && (hipLaunchKernelStruct_.bf2 == 1)
                     && (hipLaunchKernelStruct_.l1 == 1.0)
                     && (hipLaunchKernelStruct_.bf3 == 1) );
}

// Passing empty struct to a hipLaunchKernelGGL(),
// check the size of 1Byte, set  result_d4 to true if condition met
__global__ void hipLaunchKernelStructFunc4(
                    hipLaunchKernelStruct_t4 hipLaunchKernelStruct_,
                    bool* result_d4) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d4[x] =  (sizeof(hipLaunchKernelStruct_) == 1);
}

// Passing struct with pointer object to a hipLaunchKernelGGL()
__global__ void hipLaunchKernelStructFunc5(
                    hipLaunchKernelStruct_t5 hipLaunchKernelStruct_,
                    bool* result_d5) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d5[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (*hipLaunchKernelStruct_.cp == 'p'));
}

// Passing struct which is aligned to 8Byte to a hipLaunchKernelGGL(),
// set the result_d6 to true if condition met
__global__ void hipLaunchKernelStructFunc6(
                    hipLaunchKernelStruct_t6 hipLaunchKernelStruct_,
                    bool* result_d6) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    // get the address of the struct
    // size_t(p)%8 will be 0 if aligned to 8Byte address space
    int *p = (int*)(&hipLaunchKernelStruct_);
    result_d6[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (hipLaunchKernelStruct_.si == 1)
                      && ((size_t(p))%8 ==0));
}

// Passing struct which is aligned to 16Byte,
// set the result_d7 to true if condition met
__global__ void hipLaunchKernelStructFunc7(
                    hipLaunchKernelStruct_t7 hipLaunchKernelStruct_,
                    bool* result_d7) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    // get the address of the struct
    // size_t(p)%16 will be 0 if aligned to 16Byte address space
    int *p = (int*)(&hipLaunchKernelStruct_);
    result_d7[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (hipLaunchKernelStruct_.si == 1)
                      && ((size_t(p))%16 ==0) );
}

// Passing struct which is packed & aligned to 4Byte,
// set the result_d8 to true if condition met
__global__ void hipLaunchKernelStructFunc8(
                    hipLaunchKernelStruct_t8 hipLaunchKernelStruct_,
                    bool* result_d8) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    // get the address of the xth element, struct[x],
    // size_t(p)%4 will be 0 if aligned to 4Byte address space
    int *p = (int*)(&hipLaunchKernelStruct_);
    result_d8[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (hipLaunchKernelStruct_.si == 1)
                      && ((size_t(p))%4 ==0)
                      && (sizeof(hipLaunchKernelStruct_) == 4));
}

// Passing struct which is packed only, as Sam suggested, should be 4Bytes
// set the result_d8A to true if condition met
__global__ void hipLaunchKernelStructFunc8A(
                    hipLaunchKernelStruct_t8A hipLaunchKernelStruct_,
                    bool* result_d8A) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    // this is packed struct
    // the address will not be aglined in this case hence condition removed
    // only sizeof(hipLaunchKernelStruct_) will be valided
    result_d8A[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (hipLaunchKernelStruct_.si == 1)
                      && (sizeof(hipLaunchKernelStruct_) == 4));
}

// Passing struct which is aligned(4) only, as Sam suggested
// , size should be 8Bytes, set the result_d8B to true if condition met
__global__ void hipLaunchKernelStructFunc8B(
                    hipLaunchKernelStruct_t8B hipLaunchKernelStruct_,
                    bool* result_d8B) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    // get the address of the xth element, struct[x],
    // size_t(p)%4 will be 0 if aligned to 4Byte address space
    int *p = (int*)(&hipLaunchKernelStruct_);
    result_d8B[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (hipLaunchKernelStruct_.si == 1)
                      && ((size_t(p))%8 == 0)
                      && (sizeof(hipLaunchKernelStruct_) == 8));
}

// Passing struct with uint pointer object to a hipLaunchKernelGGL()
__global__ void hipLaunchKernelStructFunc9(
                    const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_,
                    bool* result_d9) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d9[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (*hipLaunchKernelStruct_.ip == 1));
}

// Passing struct with stdint types object, uintN_t, to a hipLaunchKernelGGL()
__global__ void hipLaunchKernelStructFunc10(
                    hipLaunchKernelStruct_t10 hipLaunchKernelStruct_,
                    bool* result_d10) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d10[x] =  ((hipLaunchKernelStruct_.u64 == UINT64_MAX)
                      && (hipLaunchKernelStruct_.u32 == 1)
                      && (hipLaunchKernelStruct_.u8 == UINT8_MAX));
}

// Passing struct with volatile member, to a hipLaunchKernelGGL()
__global__ void hipLaunchKernelStructFunc11(
                    hipLaunchKernelStruct_t11 hipLaunchKernelStruct_,
                    bool* result_d11) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d11[x] =  ((hipLaunchKernelStruct_.i1 == 1)
                      && (hipLaunchKernelStruct_.vint == 0));
}

// Passing struct with simple class obj, to a hipLaunchKernelGGL()
__global__ void hipLaunchKernelStructFunc12(
                    hipLaunchKernelStruct_t12 hipLaunchKernelStruct_,
                    bool* result_d12) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d12[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                      && (hipLaunchKernelStruct_.b.i == 0));
}

// Passing struct with simple __device__ func(), to a hipLaunchKernelGGL()
__global__ void hipLaunchKernelStructFunc13(
                    hipLaunchKernelStruct_t13 hipLaunchKernelStruct_,
                    bool* result_d13) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d13[x] =  ((hipLaunchKernelStruct_.i1 == 1)
                      && (hipLaunchKernelStruct_.getvalue() == 1));
}

// Passing struct with array variable, write to from device
__global__ void hipLaunchKernelStructFunc14(
                    hipLaunchKernelStruct_t14 hipLaunchKernelStruct_,
                    bool* result_d14) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    hipLaunchKernelStruct_.writeint[x] = 1;
    // set the result to true if the condition met
    result_d14[x] =  ((hipLaunchKernelStruct_.readint == 1)
                      && (hipLaunchKernelStruct_.writeint[x] == 1));
}

// Passing struct with struct with dynamic memory, new int
// the heap memory will be accessed from device
__global__ void hipLaunchKernelStructFunc15(
                    hipLaunchKernelStruct_t15 hipLaunchKernelStruct_,
                    bool* result_d15) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d15[x] =  ((hipLaunchKernelStruct_.c1 == 'c')
                       && (hipLaunchKernelStruct_.heapmem[x] == 1));
}

// Passing simple template struct
__global__ void hipLaunchKernelStructFunc16(
                    hipLaunchKernelStruct_t16<char> hipLaunchKernelStruct_,
                    bool* result_d16) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d16[x] =  (hipLaunchKernelStruct_.t1 == 'c');
}

// Passing simple explicit template struct
__global__ void hipLaunchKernelStructFunc17(
                    hipLaunchKernelStruct_t17<int> hipLaunchKernelStruct_,
                    bool* result_d17) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // set the result to true if the condition met
    result_d17[x] =  (hipLaunchKernelStruct_.t1 == 1);
}

// Passing struct and write to struct memory using __device__ func()
__global__ void hipLaunchKernelStructFunc18(
                    hipLaunchKernelStruct_t18 hipLaunchKernelStruct_,
                    bool* result_d18) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    hipLaunchKernelStruct_.setChar('c');
    // set the result to true if the condition met
    result_d18[x] =  (hipLaunchKernelStruct_.getChar() == 'c');
}

// Passing simple user defined stack implemenration,  using __device__ func()
__global__ void hipLaunchKernelStructFunc19(
                    hipLaunchKernelStruct_t19 hipLaunchKernelStruct_) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // stack should be empty after the kernel execustion, verify on host side
    hipLaunchKernelStruct_.popMe();
}

// Passing out of order initalized struct, access in-order
__global__ void hipLaunchKernelStructFunc20(
                    hipLaunchKernelStruct_t20 hipLaunchKernelStruct_,
                    bool* result_d20) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // accessing struct members in order
    result_d20[x] = (hipLaunchKernelStruct_.name == 'A'
    // strcmp(hipLaunchKernelStruct_.name, "AMD") -> strcmp is not broken
                     && hipLaunchKernelStruct_.age == 42
                     && hipLaunchKernelStruct_.rank == 2);
}

// Passing struct with bit fields
__global__ void hipLaunchKernelStructFunc21(
                    hipLaunchKernelStruct_t21 hipLaunchKernelStruct_,
                    bool* result_d21) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;

    // accessing struct members in order
    result_d21[x] = (hipLaunchKernelStruct_.i == 2
                     && hipLaunchKernelStruct_.j == 0
                     && (sizeof(hipLaunchKernelStruct_) == 1));
}

__global__ void vAdd(float* a) {}

//---
// Some wrapper macro for testing:
#define WRAP(...) __VA_ARGS__

#define MY_LAUNCH_MACRO(cmd, elapsed, quiet)                         \
    do {                                                            \
        hipDeviceSynchronize();                                     \
        cmd;                                                        \
        hipDeviceSynchronize();                                     \
    } while (0);

#define MY_LAUNCH(command, doTrace, msg)                            \
    {                                                               \
        if (doTrace) printf("TRACE: %s %s\n", msg, #command);       \
        command;                                                    \
    }

#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg)                 \
    {                                                               \
        if (doTrace) printf("TRACE: %s %s\n", msg, #command);       \
        (command);                                                  \
    }

int main() {
    // Validating memory & initial value, for result_d, result_h
    HIPASSERT(hipMallocError == hipSuccess);
    HIPASSERT(hipHostMallocError == hipSuccess);
    HIPASSERT(hipMemsetError == hipSuccess);

    // Test: Passing Struct type,  check access from device.
    ResetValidationMem();
    hipLaunchKernelStruct_t1 hipLaunchKernelStruct_h1;
    hipLaunchKernelStruct_h1.li = 1;
    hipLaunchKernelStruct_h1.lf = 1.0;
    hipLaunchKernelStruct_h1.result = false;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc1),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h1,
                    result_d);
    ResultValidation();

    // Test: Passing Struct type, checks padding
    ResetValidationMem();
    hipLaunchKernelStruct_t2 hipLaunchKernelStruct_h2;
    hipLaunchKernelStruct_h2.c1 = 'a';
    hipLaunchKernelStruct_h2.l1 = 1.0;
    hipLaunchKernelStruct_h2.c2 = 'b';
    hipLaunchKernelStruct_h2.l2 = 2.0;
    hipLaunchKernelStruct_h2.result = false;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc2),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h2,
                    result_d);
    ResultValidation();

    // Test: Passing Struct type, checks padding, assigning integer to a char
    ResetValidationMem();
    hipLaunchKernelStruct_t3 hipLaunchKernelStruct_h3;
    hipLaunchKernelStruct_h3.bf1 = 1;
    hipLaunchKernelStruct_h3.bf2 = 1;
    hipLaunchKernelStruct_h3.l1 = 1.0;
    hipLaunchKernelStruct_h3.bf3 = 1;
    hipLaunchKernelStruct_h3.result = false;
                // initialize to false, will be set to
                // true if the struct size is 1Byte, from device size
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc3),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h3,
                    result_d);
    ResultValidation();

    // Test: Passing empty struct
    ResetValidationMem();
    hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc4),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h4,
                    result_d);
    ResultValidation();

    // Test: Passing struct with pointer object to a hipLaunchKernelGGL()
    ResetValidationMem();
    hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5;
    char* cp_d5;  // This is passed as pointer to struct member
    // allocating memory for char pointer on device
    HIPCHECK(hipMalloc((void**)&cp_d5, sizeof(char)));
    HIPCHECK(hipMemset(cp_d5, 'p', sizeof(char)));
    hipLaunchKernelStruct_h5.c1 = 'c';
    hipLaunchKernelStruct_h5.cp = cp_d5;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc5),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h5,
                    result_d);
    ResultValidation();

    // Test: Passing struct with aligned(8)
    ResetValidationMem();
    hipLaunchKernelStruct_t6 hipLaunchKernelStruct_h6;
    hipLaunchKernelStruct_h6.c1 = 'c';
    hipLaunchKernelStruct_h6.si = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc6),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h6,
                    result_d);
    // alignment is broken hence disabled the validation part
    #if ENABLE_ALIGNMENT_TEST_SMALL_BAR
    ResultValidation();
    #endif

    // Test: Passing struct with aligned(16)
    ResetValidationMem();
    hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7;
    hipLaunchKernelStruct_h7.c1 = 'c';
    hipLaunchKernelStruct_h7.si = 1;
    #if ENABLE_ALIGNMENT_TEST_SMALL_BAR  // This is broken on small bar
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc7),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h7,
                    result_d);
    ResultValidation();
    #endif

    // Test: Passing struct with packed aligned to 4Bytes
    ResetValidationMem();
    hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8;
    hipLaunchKernelStruct_h8.c1 = 'c';
    hipLaunchKernelStruct_h8.si = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h8,
                    result_d);
    // packed member broken on large and small bar setup.
    #if ENABLE_PACKED_TEST
    ResultValidation();
    #endif

    // Test: Passing struct with packed to 4Bytes
    ResetValidationMem();
    hipLaunchKernelStruct_t8A hipLaunchKernelStruct_h8A;
    hipLaunchKernelStruct_h8A.c1 = 'c';
    hipLaunchKernelStruct_h8A.si = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8A),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h8A,
                    result_d);
    // packed member broken on large and small bar setup.
    #if ENABLE_PACKED_TEST
    ResultValidation();
    #endif

    // Test: Passing struct with aligned(4) to 4Bytes, size is 8Bytes
    ResetValidationMem();
    hipLaunchKernelStruct_t8B hipLaunchKernelStruct_h8B;
    hipLaunchKernelStruct_h8B.c1 = 'c';
    hipLaunchKernelStruct_h8B.si = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8B),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h8B,
                    result_d);
    // alignment is broken hence disabled the validation part
    #if ENABLE_ALIGNMENT_TEST_SMALL_BAR
    ResultValidation();
    #endif

    // Test: Passing const struct object to a hipLaunchKernelGGL()
    ResetValidationMem();
    uint32_t* ip_d9;
    // allocating memory for char pointer on device
    HIPCHECK(hipMalloc((void**)&ip_d9, sizeof(uint32_t)));
    HIPCHECK(hipMemset(ip_d9, 1, sizeof(uint32_t)));
    // ip_d9 passed as pointer to struct member, struct.ip = &ip_d9
    const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_h9 = {'c', ip_d9};
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc9),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h9,
                    result_d);
    #if ENABLE_DECLARE_INITIALIZATION_POINTER
    ResultValidation();
    #endif

    // Test: Passing struct with uintN_t as member variables
    ResetValidationMem();
    hipLaunchKernelStruct_t10 hipLaunchKernelStruct_h10;
    hipLaunchKernelStruct_h10.u64 = UINT64_MAX;
    hipLaunchKernelStruct_h10.u32 = 1;
    hipLaunchKernelStruct_h10.u8 = UINT8_MAX;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc10),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h10,
                    result_d);
    ResultValidation();

    // Test: Passing struct with uintN_t as member variables
    ResetValidationMem();
    hipLaunchKernelStruct_t11 hipLaunchKernelStruct_h11;
    hipLaunchKernelStruct_h11.i1 = 1;
    hipLaunchKernelStruct_h11.vint = 0;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc11),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h11,
                    result_d);
    ResultValidation();

    // Test: Passing struct with simple class object
    ResetValidationMem();
    hipLaunchKernelStruct_t12 hipLaunchKernelStruct_h12;
    hipLaunchKernelStruct_h12.c1 = 'c';
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc12),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h12,
                    result_d);
    #if ENABLE_CLASS_OBJ_ACCESS  // access class obj from device broken
    // Validation part of the struct, hipLaunchKernelStructFunc12
    ResultValidation();
    #endif

    // Test: Passing struct with simple __device__ func()
    ResetValidationMem();
    hipLaunchKernelStruct_t13 hipLaunchKernelStruct_h13;
    hipLaunchKernelStruct_h13.i1 = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc13),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h13,
                    result_d);
    ResultValidation();

    // Test: Passing struct with array variable, write to from device
    ResetValidationMem();
    hipLaunchKernelStruct_t14 hipLaunchKernelStruct_h14;
    hipLaunchKernelStruct_h14.readint = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc14),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h14,
                    result_d);
    ResultValidation();

    // Test: Passing struct with heap memory, read to from device
    ResetValidationMem();
    hipLaunchKernelStruct_t15 hipLaunchKernelStruct_h15;
    hipLaunchKernelStruct_h15.c1 = 'c';

    #if ENABLE_HEAP_MEMORY_ACCESS  // causing page fault here,
                                   // on small bar set
    HIPCHECK(hipMalloc(&hipLaunchKernelStruct_h15.heapmem,
                       BLOCK_DIM_SIZE*sizeof(int)));
    HIPCHECK(hipMemset(&hipLaunchKernelStruct_h15.heapmem,
                       0, BLOCK_DIM_SIZE));
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc15),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h15,
                    result_d);
    ResultValidation();
    #endif

    // Test: Passing simple template struct
    ResetValidationMem();
    hipLaunchKernelStruct_t16<char> hipLaunchKernelStruct_h16;
    hipLaunchKernelStruct_h16.t1 = 'c';
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc16),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h16,
                    result_d);
    ResultValidation();

    // Test: Passing simple explicit template struct
    ResetValidationMem();
    hipLaunchKernelStruct_t17<int> hipLaunchKernelStruct_h17;
    hipLaunchKernelStruct_h17.t1 = 1;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc17),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h17,
                    result_d);
    ResultValidation();

    // Test: Passing struct with simple __device__ func() to struct memory
    ResetValidationMem();
    hipLaunchKernelStruct_t18 hipLaunchKernelStruct_h18;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc18),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h18,
                    result_d);
    ResultValidation();

    // Test: Passing user defined stack,
    ResetValidationMem();
    hipLaunchKernelStruct_t19 hipLaunchKernelStruct_h19;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc19),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h19);
    #if ENABLE_USER_STL
    // Validation part of the struct, hipLaunchKernelStructFunc19
    HIPASSERT(hipLaunchKernelStruct_h19.stackSize() == 0);
    #endif

    // Test: Passing struct which is initiazed out of order
    // accessing same elements in order from device
    ResetValidationMem();
    hipLaunchKernelStruct_t20 hipLaunchKernelStruct_h20;
    hipLaunchKernelStruct_h20.name = 'A';
    hipLaunchKernelStruct_h20.rank = 2;
    hipLaunchKernelStruct_h20.age = 42;
    bool *result_d20, *result_h20;
    #if ENABLE_OUT_OF_ORDER_INITIALIZATION
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc20),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h20, result_d);
    ResultValidation();
    #endif

    // Test: Passing struct with bit fields operation
    // accessing same elements in order from device
    ResetValidationMem();
    hipLaunchKernelStruct_t21 hipLaunchKernelStruct_h21 =
    // out of order initalization
                     {2,0};
    bool *result_d21, *result_h21;
    hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc21),
                    dim3(BLOCK_DIM_SIZE),
                    dim3(1), 0, 0, hipLaunchKernelStruct_h21, result_d);
    #if ENABLE_BIT_FIELDS
    ResultValidation();
    #endif

    // Test: Passing the different hipLaunchParm options:
    float* Ad;
    hipMalloc((void**)&Ad, 1024);
    hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), size_t(1024), 1, 0, 0, Ad);
    hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), 1024, dim3(1), 0, 0, Ad);
    hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad);
    hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad);

    // Test: Passing hipLaunchKernelGGL inside another macro:
    float e0;
    MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024),
                   dim3(1), 0, 0, Ad), e0, j);
    MY_LAUNCH_MACRO(WRAP(hipLaunchKernelGGL(vAdd, dim3(1024),
                   dim3(1), 0, 0, Ad)), e0, j);

#ifdef EXTRA_PARENS_1
    // Don't wrap hipLaunchKernelGGL in extra set of parens:
    MY_LAUNCH_MACRO((hipLaunchKernelGGL(vAdd, dim3(1024),
                    dim3(1), 0, 0, Ad)), e0, j);
#endif

    MY_LAUNCH(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1),
              0, 0, Ad), true, "firstCall");

    float* A;
    float e1;
    MY_LAUNCH_WITH_PAREN(hipMalloc(&A, 100), true, "launch2");

#ifdef EXTRA_PARENS_2
    // MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues.
    MY_LAUNCH_WITH_PAREN(hipLaunchKernelGGL(vAdd, dim3(1024),
                         dim3(1), 0, 0, Ad), true, "firstCall");
#endif

    HIPCHECK(hipHostFree(result_h));
    HIPCHECK(hipFree(result_d));

    passed();
}


text:
package gob

// TODO(rsc): When garbage collector changes, revisit
// the allocations in this file that use unsafe.Pointer.

import (
	"bytes"
	"errors"
	"io"
	"math"
	"reflect"
	"unsafe"
)

var (
	errBadUint = errors.New("gob: encoded unsigned integer out of range")
	errBadType = errors.New("gob: unknown type id or corrupted data")
	errRange   = errors.New("gob: bad data: field numbers out of bounds")
)

// decoderState is the execution state of an instance of the decoder. A new state
// is created for nested objects.
type decoderState struct {
	dec *Decoder
	// The buffer is stored with an extra indirection because it may be replaced
	// if we load a type during decode (when reading an interface value).
	b        *bytes.Buffer
	fieldnum int // the last field number read.
	buf      []byte
	next     *decoderState // for free list
}

// We pass the bytes.Buffer separately for easier testing of the infrastructure
// without requiring a full Decoder.
func (dec *Decoder) newDecoderState(buf *bytes.Buffer) *decoderState {
	d := dec.freeList
	if d == nil {
		d = new(decoderState)
		d.dec = dec
		d.buf = make([]byte, uint64Size)
	} else {
		dec.freeList = d.next
	}
	d.b = buf
	return d
}

func (dec *Decoder) freeDecoderState(d *decoderState) {
	d.next = dec.freeList
	dec.freeList = d
}

func overflow(name string) error {
	return errors.New(`value for "` + name + `" out of range`)
}

// decodeUintReader reads an encoded unsigned integer from an io.Reader.
// Used only by the Decoder to read the message length.
func decodeUintReader(r io.Reader, buf []byte) (x uint64, width int, err error) {
	width = 1
	n, err := io.ReadFull(r, buf[0:width])
	if n == 0 {
		return
	}
	b := buf[0]
	if b <= 0x7f {
		return uint64(b), width, nil
	}
	n = -int(int8(b))
	if n > uint64Size {
		err = errBadUint
		return
	}
	width, err = io.ReadFull(r, buf[0:n])
	if err != nil {
		if err == io.EOF {
			err = io.ErrUnexpectedEOF
		}
		return
	}
	// Could check that the high byte is zero but it's not worth it.
	for _, b := range buf[0:width] {
		x = x<<8 | uint64(b)
	}
	width++ // +1 for length byte
	return
}

// decodeUint reads an encoded unsigned integer from state.r.
// Does not check for overflow.
func (state *decoderState) decodeUint() (x uint64) {
	b, err := state.b.ReadByte()
	if err != nil {
		error_(err)
	}
	if b <= 0x7f {
		return uint64(b)
	}
	n := -int(int8(b))
	if n > uint64Size {
		error_(errBadUint)
	}
	width, err := state.b.Read(state.buf[0:n])
	if err != nil {
		error_(err)
	}
	// Don't need to check error; it's safe to loop regardless.
	// Could check that the high byte is zero but it's not worth it.
	for _, b := range state.buf[0:width] {
		x = x<<8 | uint64(b)
	}
	return x
}

// decodeInt reads an encoded signed integer from state.r.
// Does not check for overflow.
func (state *decoderState) decodeInt() int64 {
	x := state.decodeUint()
	if x&1 != 0 {
		return ^int64(x >> 1)
	}
	return int64(x >> 1)
}

// decOp is the signature of a decoding operator for a given type.
type decOp func(i *decInstr, state *decoderState, p unsafe.Pointer)

// The 'instructions' of the decoding machine
type decInstr struct {
	op     decOp
	field  int     // field number of the wire type
	indir  int     // how many pointer indirections to reach the value in the struct
	offset uintptr // offset in the structure of the field to encode
	ovfl   error   // error message for overflow/underflow (for arrays, of the elements)
}

// Since the encoder writes no zeros, if we arrive at a decoder we have
// a value to extract and store.  The field number has already been read
// (it's how we knew to call this decoder).
// Each decoder is responsible for handling any indirections associated
// with the data structure.  If any pointer so reached is nil, allocation must
// be done.

// Walk the pointer hierarchy, allocating if we find a nil.  Stop one before the end.
func decIndirect(p unsafe.Pointer, indir int) unsafe.Pointer {
	for ; indir > 1; indir-- {
		if *(*unsafe.Pointer)(p) == nil {
			// Allocation required
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(unsafe.Pointer))
		}
		p = *(*unsafe.Pointer)(p)
	}
	return p
}

// ignoreUint discards a uint value with no destination.
func ignoreUint(i *decInstr, state *decoderState, p unsafe.Pointer) {
	state.decodeUint()
}

// ignoreTwoUints discards a uint value with no destination. It's used to skip
// complex values.
func ignoreTwoUints(i *decInstr, state *decoderState, p unsafe.Pointer) {
	state.decodeUint()
	state.decodeUint()
}

// decBool decodes a uint and stores it as a boolean through p.
func decBool(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(bool))
		}
		p = *(*unsafe.Pointer)(p)
	}
	*(*bool)(p) = state.decodeUint() != 0
}

// decInt8 decodes an integer and stores it as an int8 through p.
func decInt8(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(int8))
		}
		p = *(*unsafe.Pointer)(p)
	}
	v := state.decodeInt()
	if v < math.MinInt8 || math.MaxInt8 < v {
		error_(i.ovfl)
	} else {
		*(*int8)(p) = int8(v)
	}
}

// decUint8 decodes an unsigned integer and stores it as a uint8 through p.
func decUint8(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(uint8))
		}
		p = *(*unsafe.Pointer)(p)
	}
	v := state.decodeUint()
	if math.MaxUint8 < v {
		error_(i.ovfl)
	} else {
		*(*uint8)(p) = uint8(v)
	}
}

// decInt16 decodes an integer and stores it as an int16 through p.
func decInt16(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(int16))
		}
		p = *(*unsafe.Pointer)(p)
	}
	v := state.decodeInt()
	if v < math.MinInt16 || math.MaxInt16 < v {
		error_(i.ovfl)
	} else {
		*(*int16)(p) = int16(v)
	}
}

// decUint16 decodes an unsigned integer and stores it as a uint16 through p.
func decUint16(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(uint16))
		}
		p = *(*unsafe.Pointer)(p)
	}
	v := state.decodeUint()
	if math.MaxUint16 < v {
		error_(i.ovfl)
	} else {
		*(*uint16)(p) = uint16(v)
	}
}

// decInt32 decodes an integer and stores it as an int32 through p.
func decInt32(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(int32))
		}
		p = *(*unsafe.Pointer)(p)
	}
	v := state.decodeInt()
	if v < math.MinInt32 || math.MaxInt32 < v {
		error_(i.ovfl)
	} else {
		*(*int32)(p) = int32(v)
	}
}

// decUint32 decodes an unsigned integer and stores it as a uint32 through p.
func decUint32(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(uint32))
		}
		p = *(*unsafe.Pointer)(p)
	}
	v := state.decodeUint()
	if math.MaxUint32 < v {
		error_(i.ovfl)
	} else {
		*(*uint32)(p) = uint32(v)
	}
}

// decInt64 decodes an integer and stores it as an int64 through p.
func decInt64(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(int64))
		}
		p = *(*unsafe.Pointer)(p)
	}
	*(*int64)(p) = int64(state.decodeInt())
}

// decUint64 decodes an unsigned integer and stores it as a uint64 through p.
func decUint64(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(uint64))
		}
		p = *(*unsafe.Pointer)(p)
	}
	*(*uint64)(p) = uint64(state.decodeUint())
}

// Floating-point numbers are transmitted as uint64s holding the bits
// of the underlying representation.  They are sent byte-reversed, with
// the exponent end coming out first, so integer floating point numbers
// (for example) transmit more compactly.  This routine does the
// unswizzling.
func floatFromBits(u uint64) float64 {
	var v uint64
	for i := 0; i < 8; i++ {
		v <<= 8
		v |= u & 0xFF
		u >>= 8
	}
	return math.Float64frombits(v)
}

// storeFloat32 decodes an unsigned integer, treats it as a 32-bit floating-point
// number, and stores it through p. It's a helper function for float32 and complex64.
func storeFloat32(i *decInstr, state *decoderState, p unsafe.Pointer) {
	v := floatFromBits(state.decodeUint())
	av := v
	if av < 0 {
		av = -av
	}
	// +Inf is OK in both 32- and 64-bit floats.  Underflow is always OK.
	if math.MaxFloat32 < av && av <= math.MaxFloat64 {
		error_(i.ovfl)
	} else {
		*(*float32)(p) = float32(v)
	}
}

// decFloat32 decodes an unsigned integer, treats it as a 32-bit floating-point
// number, and stores it through p.
func decFloat32(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(float32))
		}
		p = *(*unsafe.Pointer)(p)
	}
	storeFloat32(i, state, p)
}

// decFloat64 decodes an unsigned integer, treats it as a 64-bit floating-point
// number, and stores it through p.
func decFloat64(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(float64))
		}
		p = *(*unsafe.Pointer)(p)
	}
	*(*float64)(p) = floatFromBits(uint64(state.decodeUint()))
}

// decComplex64 decodes a pair of unsigned integers, treats them as a
// pair of floating point numbers, and stores them as a complex64 through p.
// The real part comes first.
func decComplex64(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(complex64))
		}
		p = *(*unsafe.Pointer)(p)
	}
	storeFloat32(i, state, p)
	storeFloat32(i, state, unsafe.Pointer(uintptr(p)+unsafe.Sizeof(float32(0))))
}

// decComplex128 decodes a pair of unsigned integers, treats them as a
// pair of floating point numbers, and stores them as a complex128 through p.
// The real part comes first.
func decComplex128(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(complex128))
		}
		p = *(*unsafe.Pointer)(p)
	}
	real := floatFromBits(uint64(state.decodeUint()))
	imag := floatFromBits(uint64(state.decodeUint()))
	*(*complex128)(p) = complex(real, imag)
}

// decUint8Slice decodes a byte slice and stores through p a slice header
// describing the data.
// uint8 slices are encoded as an unsigned count followed by the raw bytes.
func decUint8Slice(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new([]uint8))
		}
		p = *(*unsafe.Pointer)(p)
	}
	n := state.decodeUint()
	if n > uint64(state.b.Len()) {
		errorf("length of []byte exceeds input size (%d bytes)", n)
	}
	slice := (*[]uint8)(p)
	if uint64(cap(*slice)) < n {
		*slice = make([]uint8, n)
	} else {
		*slice = (*slice)[0:n]
	}
	if _, err := state.b.Read(*slice); err != nil {
		errorf("error decoding []byte: %s", err)
	}
}

// decString decodes byte array and stores through p a string header
// describing the data.
// Strings are encoded as an unsigned count followed by the raw bytes.
func decString(i *decInstr, state *decoderState, p unsafe.Pointer) {
	if i.indir > 0 {
		if *(*unsafe.Pointer)(p) == nil {
			*(*unsafe.Pointer)(p) = unsafe.Pointer(new(string))
		}
		p = *(*unsafe.Pointer)(p)
	}
	n := state.decodeUint()
	if n > uint64(state.b.Len()) {
		errorf("string length exceeds input size (%d bytes)", n)
	}
	b := make([]byte, n)
	state.b.Read(b)
	// It would be a shame to do the obvious thing here,
	//	*(*string)(p) = string(b)
	// because we've already allocated the storage and this would
	// allocate again and copy.  So we do this ugly hack, which is even
	// even more unsafe than it looks as it depends the memory
	// representation of a string matching the beginning of the memory
	// representation of a byte slice (a byte slice is longer).
	*(*string)(p) = *(*string)(unsafe.Pointer(&b))
}

// ignoreUint8Array skips over the data for a byte slice value with no destination.
func ignoreUint8Array(i *decInstr, state *decoderState, p unsafe.Pointer) {
	b := make([]byte, state.decodeUint())
	state.b.Read(b)
}

// Execution engine

// The encoder engine is an array of instructions indexed by field number of the incoming
// decoder.  It is executed with random access according to field number.
type decEngine struct {
	instr    []decInstr
	numInstr int // the number of active instructions
}

// allocate makes sure storage is available for an object of underlying type rtyp
// that is indir levels of indirection through p.
func allocate(rtyp reflect.Type, p unsafe.Pointer, indir int) unsafe.Pointer {
	if indir == 0 {
		return p
	}
	up := p
	if indir > 1 {
		up = decIndirect(up, indir)
	}
	if *(*unsafe.Pointer)(up) == nil {
		// Allocate object.
		*(*unsafe.Pointer)(up) = unsafe.Pointer(reflect.New(rtyp).Pointer())
	}
	return *(*unsafe.Pointer)(up)
}

// decodeSingle decodes a top-level value that is not a struct and stores it through p.
// Such values are preceded by a zero, making them have the memory layout of a
// struct field (although with an illegal field number).
func (dec *Decoder) decodeSingle(engine *decEngine, ut *userTypeInfo, basep unsafe.Pointer) {
	state := dec.newDecoderState(&dec.buf)
	state.fieldnum = singletonField
	delta := int(state.decodeUint())
	if delta != 0 {
		errorf("decode: corrupted data: non-zero delta for singleton")
	}
	instr := &engine.instr[singletonField]
	if instr.indir != ut.indir {
		errorf("internal error: inconsistent indirection instr %d ut %d", instr.indir, ut.indir)
	}
	ptr := basep // offset will be zero
	if instr.indir > 1 {
		ptr = decIndirect(ptr, instr.indir)
	}
	instr.op(instr, state, ptr)
	dec.freeDecoderState(state)
}

// decodeStruct decodes a top-level struct and stores it through p.
// Indir is for the value, not the type.  At the time of the call it may
// differ from ut.indir, which was computed when the engine was built.
// This state cannot arise for decodeSingle, which is called directly
// from the user's value, not from the innards of an engine.
func (dec *Decoder) decodeStruct(engine *decEngine, ut *userTypeInfo, p unsafe.Pointer, indir int) {
	p = allocate(ut.base, p, indir)
	state := dec.newDecoderState(&dec.buf)
	state.fieldnum = -1
	basep := p
	for state.b.Len() > 0 {
		delta := int(state.decodeUint())
		if delta < 0 {
			errorf("decode: corrupted data: negative delta")
		}
		if delta == 0 { // struct terminator is zero delta fieldnum
			break
		}
		fieldnum := state.fieldnum + delta
		if fieldnum >= len(engine.instr) {
			error_(errRange)
			break
		}
		instr := &engine.instr[fieldnum]
		p := unsafe.Pointer(uintptr(basep) + instr.offset)
		if instr.indir > 1 {
			p = decIndirect(p, instr.indir)
		}
		instr.op(instr, state, p)
		state.fieldnum = fieldnum
	}
	dec.freeDecoderState(state)
}

// ignoreStruct discards the data for a struct with no destination.
func (dec *Decoder) ignoreStruct(engine *decEngine) {
	state := dec.newDecoderState(&dec.buf)
	state.fieldnum = -1
	for state.b.Len() > 0 {
		delta := int(state.decodeUint())
		if delta < 0 {
			errorf("ignore decode: corrupted data: negative delta")
		}
		if delta == 0 { // struct terminator is zero delta fieldnum
			break
		}
		fieldnum := state.fieldnum + delta
		if fieldnum >= len(engine.instr) {
			error_(errRange)
		}
		instr := &engine.instr[fieldnum]
		instr.op(instr, state, unsafe.Pointer(nil))
		state.fieldnum = fieldnum
	}
	dec.freeDecoderState(state)
}

// ignoreSingle discards the data for a top-level non-struct value with no
// destination. It's used when calling Decode with a nil value.
func (dec *Decoder) ignoreSingle(engine *decEngine) {
	state := dec.newDecoderState(&dec.buf)
	state.fieldnum = singletonField
	delta := int(state.decodeUint())
	if delta != 0 {
		errorf("decode: corrupted data: non-zero delta for singleton")
	}
	instr := &engine.instr[singletonField]
	instr.op(instr, state, unsafe.Pointer(nil))
	dec.freeDecoderState(state)
}

// decodeArrayHelper does the work for decoding arrays and slices.
func (dec *Decoder) decodeArrayHelper(state *decoderState, p unsafe.Pointer, elemOp decOp, elemWid uintptr, length, elemIndir int, ovfl error) {
	instr := &decInstr{elemOp, 0, elemIndir, 0, ovfl}
	for i := 0; i < length; i++ {
		if state.b.Len() == 0 {
			errorf("decoding array or slice: length exceeds input size (%d elements)", length)
		}
		up := p
		if elemIndir > 1 {
			up = decIndirect(up, elemIndir)
		}
		elemOp(instr, state, up)
		p = unsafe.Pointer(uintptr(p) + elemWid)
	}
}

// decodeArray decodes an array and stores it through p, that is, p points to the zeroth element.
// The length is an unsigned integer preceding the elements.  Even though the length is redundant
// (it's part of the type), it's a useful check and is included in the encoding.
func (dec *Decoder) decodeArray(atyp reflect.Type, state *decoderState, p unsafe.Pointer, elemOp decOp, elemWid uintptr, length, indir, elemIndir int, ovfl error) {
	if indir > 0 {
		p = allocate(atyp, p, 1) // All but the last level has been allocated by dec.Indirect
	}
	if n := state.decodeUint(); n != uint64(length) {
		errorf("length mismatch in decodeArray")
	}
	dec.decodeArrayHelper(state, p, elemOp, elemWid, length, elemIndir, ovfl)
}

// decodeIntoValue is a helper for map decoding.  Since maps are decoded using reflection,
// unlike the other items we can't use a pointer directly.
func decodeIntoValue(state *decoderState, op decOp, indir int, v reflect.Value, ovfl error) reflect.Value {
	instr := &decInstr{op, 0, indir, 0, ovfl}
	up := unsafeAddr(v)
	if indir > 1 {
		up = decIndirect(up, indir)
	}
	op(instr, state, up)
	return v
}

// decodeMap decodes a map and stores its header through p.
// Maps are encoded as a length followed by key:value pairs.
// Because the internals of maps are not visible to us, we must
// use reflection rather than pointer magic.
func (dec *Decoder) decodeMap(mtyp reflect.Type, state *decoderState, p unsafe.Pointer, keyOp, elemOp decOp, indir, keyIndir, elemIndir int, ovfl error) {
	if indir > 0 {
		p = allocate(mtyp, p, 1) // All but the last level has been allocated by dec.Indirect
	}
	up := unsafe.Pointer(p)
	if *(*unsafe.Pointer)(up) == nil { // maps are represented as a pointer in the runtime
		// Allocate map.
		*(*unsafe.Pointer)(up) = unsafe.Pointer(reflect.MakeMap(mtyp).Pointer())
	}
	// Maps cannot be accessed by moving addresses around the way
	// that slices etc. can.  We must recover a full reflection value for
	// the iteration.
	v := reflect.NewAt(mtyp, unsafe.Pointer(p)).Elem()
	n := int(state.decodeUint())
	for i := 0; i < n; i++ {
		key := decodeIntoValue(state, keyOp, keyIndir, allocValue(mtyp.Key()), ovfl)
		elem := decodeIntoValue(state, elemOp, elemIndir, allocValue(mtyp.Elem()), ovfl)
		v.SetMapIndex(key, elem)
	}
}

// ignoreArrayHelper does the work for discarding arrays and slices.
func (dec *Decoder) ignoreArrayHelper(state *decoderState, elemOp decOp, length int) {
	instr := &decInstr{elemOp, 0, 0, 0, errors.New("no error")}
	for i := 0; i < length; i++ {
		elemOp(instr, state, nil)
	}
}

// ignoreArray discards the data for an array value with no destination.
func (dec *Decoder) ignoreArray(state *decoderState, elemOp decOp, length int) {
	if n := state.decodeUint(); n != uint64(length) {
		errorf("length mismatch in ignoreArray")
	}
	dec.ignoreArrayHelper(state, elemOp, length)
}

// ignoreMap discards the data for a map value with no destination.
func (dec *Decoder) ignoreMap(state *decoderState, keyOp, elemOp decOp) {
	n := int(state.decodeUint())
	keyInstr := &decInstr{keyOp, 0, 0, 0, errors.New("no error")}
	elemInstr := &decInstr{elemOp, 0, 0, 0, errors.New("no error")}
	for i := 0; i < n; i++ {
		keyOp(keyInstr, state, nil)
		elemOp(elemInstr, state, nil)
	}
}

// decodeSlice decodes a slice and stores the slice header through p.
// Slices are encoded as an unsigned length followed by the elements.
func (dec *Decoder) decodeSlice(atyp reflect.Type, state *decoderState, p uintptr, elemOp decOp, elemWid uintptr, indir, elemIndir int, ovfl error) {
	nr := state.decodeUint()
	n := int(nr)
	if indir > 0 {
		up := unsafe.Pointer(p)
		if *(*unsafe.Pointer)(up) == nil {
			// Allocate the slice header.
			*(*unsafe.Pointer)(up) = unsafe.Pointer(new([]unsafe.Pointer))
		}
		p = *(*uintptr)(up)
	}
	// Allocate storage for the slice elements, that is, the underlying array,
	// if the existing slice does not have the capacity.
	// Always write a header at p.
	hdrp := (*reflect.SliceHeader)(unsafe.Pointer(p))
	if hdrp.Cap < n {
		hdrp.Data = reflect.MakeSlice(atyp, n, n).Pointer()
		hdrp.Cap = n
	}
	hdrp.Len = n
	dec.decodeArrayHelper(state, unsafe.Pointer(hdrp.Data), elemOp, elemWid, n, elemIndir, ovfl)
}

// ignoreSlice skips over the data for a slice value with no destination.
func (dec *Decoder) ignoreSlice(state *decoderState, elemOp decOp) {
	dec.ignoreArrayHelper(state, elemOp, int(state.decodeUint()))
}

// setInterfaceValue sets an interface value to a concrete value,
// but first it checks that the assignment will succeed.
func setInterfaceValue(ivalue reflect.Value, value reflect.Value) {
	if !value.Type().AssignableTo(ivalue.Type()) {
		errorf("cannot assign value of type %s to %s", value.Type(), ivalue.Type())
	}
	ivalue.Set(value)
}

// decodeInterface decodes an interface value and stores it through p.
// Interfaces are encoded as the name of a concrete type followed by a value.
// If the name is empty, the value is nil and no value is sent.
func (dec *Decoder) decodeInterface(ityp reflect.Type, state *decoderState, p unsafe.Pointer, indir int) {
	// Create a writable interface reflect.Value.  We need one even for the nil case.
	ivalue := allocValue(ityp)
	// Read the name of the concrete type.
	nr := state.decodeUint()
	if nr < 0 || nr > 1<<31 { // zero is permissible for anonymous types
		errorf("invalid type name length %d", nr)
	}
	b := make([]byte, nr)
	state.b.Read(b)
	name := string(b)
	if name == "" {
		// Copy the representation of the nil interface value to the target.
		// This is horribly unsafe and special.
		if indir > 0 {
			p = allocate(ityp, p, 1) // All but the last level has been allocated by dec.Indirect
		}
		*(*[2]uintptr)(unsafe.Pointer(p)) = ivalue.InterfaceData()
		return
	}
	if len(name) > 1024 {
		errorf("name too long (%d bytes): %.20q...", len(name), name)
	}
	// The concrete type must be registered.
	registerLock.RLock()
	typ, ok := nameToConcreteType[name]
	registerLock.RUnlock()
	if !ok {
		errorf("name not registered for interface: %q", name)
	}
	// Read the type id of the concrete value.
	concreteId := dec.decodeTypeSequence(true)
	if concreteId < 0 {
		error_(dec.err)
	}
	// Byte count of value is next; we don't care what it is (it's there
	// in case we want to ignore the value by skipping it completely).
	state.decodeUint()
	// Read the concrete value.
	value := allocValue(typ)
	dec.decodeValue(concreteId, value)
	if dec.err != nil {
		error_(dec.err)
	}
	// Allocate the destination interface value.
	if indir > 0 {
		p = allocate(ityp, p, 1) // All but the last level has been allocated by dec.Indirect
	}
	// Assign the concrete value to the interface.
	// Tread carefully; it might not satisfy the interface.
	setInterfaceValue(ivalue, value)
	// Copy the representation of the interface value to the target.
	// This is horribly unsafe and special.
	*(*[2]uintptr)(unsafe.Pointer(p)) = ivalue.InterfaceData()
}

// ignoreInterface discards the data for an interface value with no destination.
func (dec *Decoder) ignoreInterface(state *decoderState) {
	// Read the name of the concrete type.
	b := make([]byte, state.decodeUint())
	_, err := state.b.Read(b)
	if err != nil {
		error_(err)
	}
	id := dec.decodeTypeSequence(true)
	if id < 0 {
		error_(dec.err)
	}
	// At this point, the decoder buffer contains a delimited value. Just toss it.
	state.b.Next(int(state.decodeUint()))
}

// decodeGobDecoder decodes something implementing the GobDecoder interface.
// The data is encoded as a byte slice.
func (dec *Decoder) decodeGobDecoder(state *decoderState, v reflect.Value) {
	// Read the bytes for the value.
	b := make([]byte, state.decodeUint())
	_, err := state.b.Read(b)
	if err != nil {
		error_(err)
	}
	// We know it's a GobDecoder, so just call the method directly.
	err = v.Interface().(GobDecoder).GobDecode(b)
	if err != nil {
		error_(err)
	}
}

// ignoreGobDecoder discards the data for a GobDecoder value with no destination.
func (dec *Decoder) ignoreGobDecoder(state *decoderState) {
	// Read the bytes for the value.
	b := make([]byte, state.decodeUint())
	_, err := state.b.Read(b)
	if err != nil {
		error_(err)
	}
}

// Index by Go types.
var decOpTable = [...]decOp{
	reflect.Bool:       decBool,
	reflect.Int8:       decInt8,
	reflect.Int16:      decInt16,
	reflect.Int32:      decInt32,
	reflect.Int64:      decInt64,
	reflect.Uint8:      decUint8,
	reflect.Uint16:     decUint16,
	reflect.Uint32:     decUint32,
	reflect.Uint64:     decUint64,
	reflect.Float32:    decFloat32,
	reflect.Float64:    decFloat64,
	reflect.Complex64:  decComplex64,
	reflect.Complex128: decComplex128,
	reflect.String:     decString,
}

// Indexed by gob types.  tComplex will be added during type.init().
var decIgnoreOpMap = map[typeId]decOp{
	tBool:    ignoreUint,
	tInt:     ignoreUint,
	tUint:    ignoreUint,
	tFloat:   ignoreUint,
	tBytes:   ignoreUint8Array,
	tString:  ignoreUint8Array,
	tComplex: ignoreTwoUints,
}

// decOpFor returns the decoding op for the base type under rt and
// the indirection count to reach it.
func (dec *Decoder) decOpFor(wireId typeId, rt reflect.Type, name string, inProgress map[reflect.Type]*decOp) (*decOp, int) {
	ut := userType(rt)
	// If the type implements GobEncoder, we handle it without further processing.
	if ut.isGobDecoder {
		return dec.gobDecodeOpFor(ut)
	}
	// If this type is already in progress, it's a recursive type (e.g. map[string]*T).
	// Return the pointer to the op we're already building.
	if opPtr := inProgress[rt]; opPtr != nil {
		return opPtr, ut.indir
	}
	typ := ut.base
	indir := ut.indir
	var op decOp
	k := typ.Kind()
	if int(k) < len(decOpTable) {
		op = decOpTable[k]
	}
	if op == nil {
		inProgress[rt] = &op
		// Special cases
		switch t := typ; t.Kind() {
		case reflect.Array:
			name = "element of " + name
			elemId := dec.wireType[wireId].ArrayT.Elem
			elemOp, elemIndir := dec.decOpFor(elemId, t.Elem(), name, inProgress)
			ovfl := overflow(name)
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.decodeArray(t, state, p, *elemOp, t.Elem().Size(), t.Len(), i.indir, elemIndir, ovfl)
			}

		case reflect.Map:
			keyId := dec.wireType[wireId].MapT.Key
			elemId := dec.wireType[wireId].MapT.Elem
			keyOp, keyIndir := dec.decOpFor(keyId, t.Key(), "key of "+name, inProgress)
			elemOp, elemIndir := dec.decOpFor(elemId, t.Elem(), "element of "+name, inProgress)
			ovfl := overflow(name)
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.decodeMap(t, state, p, *keyOp, *elemOp, i.indir, keyIndir, elemIndir, ovfl)
			}

		case reflect.Slice:
			name = "element of " + name
			if t.Elem().Kind() == reflect.Uint8 {
				op = decUint8Slice
				break
			}
			var elemId typeId
			if tt, ok := builtinIdToType[wireId]; ok {
				elemId = tt.(*sliceType).Elem
			} else {
				elemId = dec.wireType[wireId].SliceT.Elem
			}
			elemOp, elemIndir := dec.decOpFor(elemId, t.Elem(), name, inProgress)
			ovfl := overflow(name)
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.decodeSlice(t, state, uintptr(p), *elemOp, t.Elem().Size(), i.indir, elemIndir, ovfl)
			}

		case reflect.Struct:
			// Generate a closure that calls out to the engine for the nested type.
			enginePtr, err := dec.getDecEnginePtr(wireId, userType(typ))
			if err != nil {
				error_(err)
			}
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				// indirect through enginePtr to delay evaluation for recursive structs.
				dec.decodeStruct(*enginePtr, userType(typ), p, i.indir)
			}
		case reflect.Interface:
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.decodeInterface(t, state, p, i.indir)
			}
		}
	}
	if op == nil {
		errorf("decode can't handle type %s", rt)
	}
	return &op, indir
}

// decIgnoreOpFor returns the decoding op for a field that has no destination.
func (dec *Decoder) decIgnoreOpFor(wireId typeId) decOp {
	op, ok := decIgnoreOpMap[wireId]
	if !ok {
		if wireId == tInterface {
			// Special case because it's a method: the ignored item might
			// define types and we need to record their state in the decoder.
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.ignoreInterface(state)
			}
			return op
		}
		// Special cases
		wire := dec.wireType[wireId]
		switch {
		case wire == nil:
			errorf("bad data: undefined type %s", wireId.string())
		case wire.ArrayT != nil:
			elemId := wire.ArrayT.Elem
			elemOp := dec.decIgnoreOpFor(elemId)
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.ignoreArray(state, elemOp, wire.ArrayT.Len)
			}

		case wire.MapT != nil:
			keyId := dec.wireType[wireId].MapT.Key
			elemId := dec.wireType[wireId].MapT.Elem
			keyOp := dec.decIgnoreOpFor(keyId)
			elemOp := dec.decIgnoreOpFor(elemId)
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.ignoreMap(state, keyOp, elemOp)
			}

		case wire.SliceT != nil:
			elemId := wire.SliceT.Elem
			elemOp := dec.decIgnoreOpFor(elemId)
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.ignoreSlice(state, elemOp)
			}

		case wire.StructT != nil:
			// Generate a closure that calls out to the engine for the nested type.
			enginePtr, err := dec.getIgnoreEnginePtr(wireId)
			if err != nil {
				error_(err)
			}
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				// indirect through enginePtr to delay evaluation for recursive structs
				state.dec.ignoreStruct(*enginePtr)
			}

		case wire.GobEncoderT != nil:
			op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
				state.dec.ignoreGobDecoder(state)
			}
		}
	}
	if op == nil {
		errorf("bad data: ignore can't handle type %s", wireId.string())
	}
	return op
}

// gobDecodeOpFor returns the op for a type that is known to implement
// GobDecoder.
func (dec *Decoder) gobDecodeOpFor(ut *userTypeInfo) (*decOp, int) {
	rcvrType := ut.user
	if ut.decIndir == -1 {
		rcvrType = reflect.PtrTo(rcvrType)
	} else if ut.decIndir > 0 {
		for i := int8(0); i < ut.decIndir; i++ {
			rcvrType = rcvrType.Elem()
		}
	}
	var op decOp
	op = func(i *decInstr, state *decoderState, p unsafe.Pointer) {
		// Caller has gotten us to within one indirection of our value.
		if i.indir > 0 {
			if *(*unsafe.Pointer)(p) == nil {
				*(*unsafe.Pointer)(p) = unsafe.Pointer(reflect.New(ut.base).Pointer())
			}
		}
		// Now p is a pointer to the base type.  Do we need to climb out to
		// get to the receiver type?
		var v reflect.Value
		if ut.decIndir == -1 {
			v = reflect.NewAt(rcvrType, unsafe.Pointer(&p)).Elem()
		} else {
			v = reflect.NewAt(rcvrType, p).Elem()
		}
		state.dec.decodeGobDecoder(state, v)
	}
	return &op, int(ut.indir)

}

// compatibleType asks: Are these two gob Types compatible?
// Answers the question for basic types, arrays, maps and slices, plus
// GobEncoder/Decoder pairs.
// Structs are considered ok; fields will be checked later.
func (dec *Decoder) compatibleType(fr reflect.Type, fw typeId, inProgress map[reflect.Type]typeId) bool {
	if rhs, ok := inProgress[fr]; ok {
		return rhs == fw
	}
	inProgress[fr] = fw
	ut := userType(fr)
	wire, ok := dec.wireType[fw]
	// If fr is a GobDecoder, the wire type must be GobEncoder.
	// And if fr is not a GobDecoder, the wire type must not be either.
	if ut.isGobDecoder != (ok && wire.GobEncoderT != nil) { // the parentheses look odd but are correct.
		return false
	}
	if ut.isGobDecoder { // This test trumps all others.
		return true
	}
	switch t := ut.base; t.Kind() {
	default:
		// chan, etc: cannot handle.
		return false
	case reflect.Bool:
		return fw == tBool
	case reflect.Int, reflect.Int8, reflect.Int16, reflect.Int32, reflect.Int64:
		return fw == tInt
	case reflect.Uint, reflect.Uint8, reflect.Uint16, reflect.Uint32, reflect.Uint64, reflect.Uintptr:
		return fw == tUint
	case reflect.Float32, reflect.Float64:
		return fw == tFloat
	case reflect.Complex64, reflect.Complex128:
		return fw == tComplex
	case reflect.String:
		return fw == tString
	case reflect.Interface:
		return fw == tInterface
	case reflect.Array:
		if !ok || wire.ArrayT == nil {
			return false
		}
		array := wire.ArrayT
		return t.Len() == array.Len && dec.compatibleType(t.Elem(), array.Elem, inProgress)
	case reflect.Map:
		if !ok || wire.MapT == nil {
			return false
		}
		MapType := wire.MapT
		return dec.compatibleType(t.Key(), MapType.Key, inProgress) && dec.compatibleType(t.Elem(), MapType.Elem, inProgress)
	case reflect.Slice:
		// Is it an array of bytes?
		if t.Elem().Kind() == reflect.Uint8 {
			return fw == tBytes
		}
		// Extract and compare element types.
		var sw *sliceType
		if tt, ok := builtinIdToType[fw]; ok {
			sw, _ = tt.(*sliceType)
		} else if wire != nil {
			sw = wire.SliceT
		}
		elem := userType(t.Elem()).base
		return sw != nil && dec.compatibleType(elem, sw.Elem, inProgress)
	case reflect.Struct:
		return true
	}
}

// typeString returns a human-readable description of the type identified by remoteId.
func (dec *Decoder) typeString(remoteId typeId) string {
	if t := idToType[remoteId]; t != nil {
		// globally known type.
		return t.string()
	}
	return dec.wireType[remoteId].string()
}

// compileSingle compiles the decoder engine for a non-struct top-level value, including
// GobDecoders.
func (dec *Decoder) compileSingle(remoteId typeId, ut *userTypeInfo) (engine *decEngine, err error) {
	rt := ut.user
	engine = new(decEngine)
	engine.instr = make([]decInstr, 1) // one item
	name := rt.String()                // best we can do
	if !dec.compatibleType(rt, remoteId, make(map[reflect.Type]typeId)) {
		remoteType := dec.typeString(remoteId)
		// Common confusing case: local interface type, remote concrete type.
		if ut.base.Kind() == reflect.Interface && remoteId != tInterface {
			return nil, errors.New("gob: local interface type " + name + " can only be decoded from remote interface type; received concrete type " + remoteType)
		}
		return nil, errors.New("gob: decoding into local type " + name + ", received remote type " + remoteType)
	}
	op, indir := dec.decOpFor(remoteId, rt, name, make(map[reflect.Type]*decOp))
	ovfl := errors.New(`value for "` + name + `" out of range`)
	engine.instr[singletonField] = decInstr{*op, singletonField, indir, 0, ovfl}
	engine.numInstr = 1
	return
}

// compileIgnoreSingle compiles the decoder engine for a non-struct top-level value that will be discarded.
func (dec *Decoder) compileIgnoreSingle(remoteId typeId) (engine *decEngine, err error) {
	engine = new(decEngine)
	engine.instr = make([]decInstr, 1) // one item
	op := dec.decIgnoreOpFor(remoteId)
	ovfl := overflow(dec.typeString(remoteId))
	engine.instr[0] = decInstr{op, 0, 0, 0, ovfl}
	engine.numInstr = 1
	return
}

// compileDec compiles the decoder engine for a value.  If the value is not a struct,
// it calls out to compileSingle.
func (dec *Decoder) compileDec(remoteId typeId, ut *userTypeInfo) (engine *decEngine, err error) {
	rt := ut.base
	srt := rt
	if srt.Kind() != reflect.Struct ||
		ut.isGobDecoder {
		return dec.compileSingle(remoteId, ut)
	}
	var wireStruct *structType
	// Builtin types can come from global pool; the rest must be defined by the decoder.
	// Also we know we're decoding a struct now, so the client must have sent one.
	if t, ok := builtinIdToType[remoteId]; ok {
		wireStruct, _ = t.(*structType)
	} else {
		wire := dec.wireType[remoteId]
		if wire == nil {
			error_(errBadType)
		}
		wireStruct = wire.StructT
	}
	if wireStruct == nil {
		errorf("type mismatch in decoder: want struct type %s; got non-struct", rt)
	}
	engine = new(decEngine)
	engine.instr = make([]decInstr, len(wireStruct.Field))
	seen := make(map[reflect.Type]*decOp)
	// Loop over the fields of the wire type.
	for fieldnum := 0; fieldnum < len(wireStruct.Field); fieldnum++ {
		wireField := wireStruct.Field[fieldnum]
		if wireField.Name == "" {
			errorf("empty name for remote field of type %s", wireStruct.Name)
		}
		ovfl := overflow(wireField.Name)
		// Find the field of the local type with the same name.
		localField, present := srt.FieldByName(wireField.Name)
		// TODO(r): anonymous names
		if !present || !isExported(wireField.Name) {
			op := dec.decIgnoreOpFor(wireField.Id)
			engine.instr[fieldnum] = decInstr{op, fieldnum, 0, 0, ovfl}
			continue
		}
		if !dec.compatibleType(localField.Type, wireField.Id, make(map[reflect.Type]typeId)) {
			errorf("wrong type (%s) for received field %s.%s", localField.Type, wireStruct.Name, wireField.Name)
		}
		op, indir := dec.decOpFor(wireField.Id, localField.Type, localField.Name, seen)
		engine.instr[fieldnum] = decInstr{*op, fieldnum, indir, uintptr(localField.Offset), ovfl}
		engine.numInstr++
	}
	return
}

// getDecEnginePtr returns the engine for the specified type.
func (dec *Decoder) getDecEnginePtr(remoteId typeId, ut *userTypeInfo) (enginePtr **decEngine, err error) {
	rt := ut.user
	decoderMap, ok := dec.decoderCache[rt]
	if !ok {
		decoderMap = make(map[typeId]**decEngine)
		dec.decoderCache[rt] = decoderMap
	}
	if enginePtr, ok = decoderMap[remoteId]; !ok {
		// To handle recursive types, mark this engine as underway before compiling.
		enginePtr = new(*decEngine)
		decoderMap[remoteId] = enginePtr
		*enginePtr, err = dec.compileDec(remoteId, ut)
		if err != nil {
			delete(decoderMap, remoteId)
		}
	}
	return
}

// emptyStruct is the type we compile into when ignoring a struct value.
type emptyStruct struct{}

var emptyStructType = reflect.TypeOf(emptyStruct{})

// getDecEnginePtr returns the engine for the specified type when the value is to be discarded.
func (dec *Decoder) getIgnoreEnginePtr(wireId typeId) (enginePtr **decEngine, err error) {
	var ok bool
	if enginePtr, ok = dec.ignorerCache[wireId]; !ok {
		// To handle recursive types, mark this engine as underway before compiling.
		enginePtr = new(*decEngine)
		dec.ignorerCache[wireId] = enginePtr
		wire := dec.wireType[wireId]
		if wire != nil && wire.StructT != nil {
			*enginePtr, err = dec.compileDec(wireId, userType(emptyStructType))
		} else {
			*enginePtr, err = dec.compileIgnoreSingle(wireId)
		}
		if err != nil {
			delete(dec.ignorerCache, wireId)
		}
	}
	return
}

// decodeValue decodes the data stream representing a value and stores it in val.
func (dec *Decoder) decodeValue(wireId typeId, val reflect.Value) {
	defer catchError(&dec.err)
	// If the value is nil, it means we should just ignore this item.
	if !val.IsValid() {
		dec.decodeIgnoredValue(wireId)
		return
	}
	// Dereference down to the underlying type.
	ut := userType(val.Type())
	base := ut.base
	var enginePtr **decEngine
	enginePtr, dec.err = dec.getDecEnginePtr(wireId, ut)
	if dec.err != nil {
		return
	}
	engine := *enginePtr
	if st := base; st.Kind() == reflect.Struct && !ut.isGobDecoder {
		if engine.numInstr == 0 && st.NumField() > 0 && len(dec.wireType[wireId].StructT.Field) > 0 {
			name := base.Name()
			errorf("type mismatch: no fields matched compiling decoder for %s", name)
		}
		dec.decodeStruct(engine, ut, unsafeAddr(val), ut.indir)
	} else {
		dec.decodeSingle(engine, ut, unsafeAddr(val))
	}
}

// decodeIgnoredValue decodes the data stream representing a value of the specified type and discards it.
func (dec *Decoder) decodeIgnoredValue(wireId typeId) {
	var enginePtr **decEngine
	enginePtr, dec.err = dec.getIgnoreEnginePtr(wireId)
	if dec.err != nil {
		return
	}
	wire := dec.wireType[wireId]
	if wire != nil && wire.StructT != nil {
		dec.ignoreStruct(*enginePtr)
	} else {
		dec.ignoreSingle(*enginePtr)
	}
}

func init() {
	var iop, uop decOp
	switch reflect.TypeOf(int(0)).Bits() {
	case 32:
		iop = decInt32
		uop = decUint32
	case 64:
		iop = decInt64
		uop = decUint64
	default:
		panic("gob: unknown size of int/uint")
	}
	decOpTable[reflect.Int] = iop
	decOpTable[reflect.Uint] = uop

	// Finally uintptr
	switch reflect.TypeOf(uintptr(0)).Bits() {
	case 32:
		uop = decUint32
	case 64:
		uop = decUint64
	default:
		panic("gob: unknown size of uintptr")
	}
	decOpTable[reflect.Uintptr] = uop
}

// Gob assumes it can call UnsafeAddr on any Value
// in order to get a pointer it can copy data from.
// Values that have just been created and do not point
// into existing structs or slices cannot be addressed,
// so simulate it by returning a pointer to a copy.
// Each call allocates once.
func unsafeAddr(v reflect.Value) unsafe.Pointer {
	if v.CanAddr() {
		return unsafe.Pointer(v.UnsafeAddr())
	}
	x := reflect.New(v.Type()).Elem()
	x.Set(v)
	return unsafe.Pointer(x.UnsafeAddr())
}

// Gob depends on being able to take the address
// of zeroed Values it creates, so use this wrapper instead
// of the standard reflect.Zero.
// Each call allocates once.
func allocValue(t reflect.Type) reflect.Value {
	return reflect.New(t).Elem()
}

redpajama_set_name: RedPajamaGithub


text:
using UnityEngine;

namespace BehaviorDesigner.Runtime.Tasks.Basic.UnityPhysics2D
{
    [TaskCategory("Basic/Physics2D")]
    [TaskDescription("Returns success if there is any collider intersecting the line between start and end")]
    [HelpURL("http://www.opsive.com/assets/BehaviorDesigner/documentation.php?id=118")]
    public class Linecast : Action
    {
        [Tooltip("The starting position of the linecast.")]
        SharedVector2 startPosition;
        [Tooltip("The ending position of the linecast.")]
        SharedVector2 endPosition;
        [Tooltip("Selectively ignore colliders.")]
        public LayerMask layerMask = -1;

        public override TaskStatus OnUpdate()
        {
            return Physics2D.Linecast(startPosition.Value, endPosition.Value, layerMask) ? TaskStatus.Success : TaskStatus.Failure;
        }

        public override void OnReset()
        {
            startPosition = Vector2.zero;
            endPosition = Vector2.zero;
            layerMask = -1;
        }
    }
}


text:
"""
Portions copyright © 2007, Stephen John Machin, Lingfo Pty Ltd
All rights reserved.

Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:

1. Redistributions of source code must retain the above copyright notice,
this list of conditions and the following disclaimer.

2. 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.

3. None of the names of Stephen John Machin, Lingfo Pty Ltd and any
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 OWNER 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.
"""

"""
  Copyright (C) 2005 Roman V. Kiseliov
  All rights reserved.

  Redistribution and use in source and binary forms, with or without
  modification, are permitted provided that the following conditions
  are met:

  1. Redistributions of source code must retain the above copyright
     notice, this list of conditions and the following disclaimer.

  2. 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.

  3. All advertising materials mentioning features or use of this
     software must display the following acknowledgment:
     "This product includes software developed by
      Roman V. Kiseliov <roman@kiseliov.ru>."

  4. Redistributions of any form whatsoever must retain the following
     acknowledgment:
     "This product includes software developed by
      Roman V. Kiseliov <roman@kiseliov.ru>."

  THIS SOFTWARE IS PROVIDED BY Roman V. Kiseliov ``AS IS'' AND ANY
  EXPRESSED 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 Roman V. Kiseliov OR
  ITS 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.

Roman V. Kiseliov
Russia
Kursk
Libknecht St., 4

+7(0712)56-09-83

<roman@kiseliov.ru>
Subject: pyExcelerator
"""

"""
Portions of xlwt.Utils based on:
pyXLWriter - A library for generating Excel Spreadsheets

The licensing of pyXLWriter is as follows:

 Copyright (c) 2004 Evgeny Filatov <fufff@users.sourceforge.net>
 Copyright (c) 2002-2004 John McNamara (Perl Spreadsheet::WriteExcel)

 This library is free software; you can redistribute it and/or modify it
 under the terms of the GNU Lesser General Public License as published by
 the Free Software Foundation; either version 2.1 of the License, or
 (at your option) any later version.

 This library 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 Lesser
 General Public License for more details:

 http://www.gnu.org/licenses/lgpl.html

 pyXLWriter also makes reference to the PERL Spreadsheet::WriteExcel as follows:

 ----------------------------------------------------------------------------
  This module was written/ported from PERL Spreadsheet::WriteExcel module
  The author of the PERL Spreadsheet::WriteExcel module is John McNamara
  <jmcnamara@cpan.org>
 ----------------------------------------------------------------------------

"""


text:
<?php declare(strict_types = 1);

/**
 * Test: Annotation\Controller\Negotiations
 */

require_once __DIR__ . '/../../../bootstrap.php';

use Apitte\Core\Annotation\Controller\Negotiation;
use Apitte\Core\Annotation\Controller\Negotiations;
use Doctrine\Common\Annotations\AnnotationException;
use Tester\Assert;

// OK
test(function (): void {
	$negotiations = new Negotiations([
		$negotiation1 = new Negotiation('json', true),
		$negotiation2 = new Negotiation('xml', false),
		$negotiation3 = new Negotiation('csv', false),
	]);

	Assert::same([$negotiation1, $negotiation2, $negotiation3], $negotiations->getNegotiations());
});

// Exception - empty negotiations
test(function (): void {
	Assert::exception(function (): void {
		new Negotiations([]);
	}, AnnotationException::class, 'Empty @Negotiations given');
});


text:
package cola.machine.game.myblocks.skill;

import cola.machine.game.myblocks.model.textture.BoneBlock;
import cola.machine.game.myblocks.model.textture.TextureInfo;

/**
 * Created by luying on 16/9/25.
 */
public class SkillDefinition {
   int id;//id
    String name;// name

    public int getDamageValue() {
        return damageValue;
    }

    public void setDamageValue(int damageValue) {
        this.damageValue = damageValue;
    }

    int damageValue;
    String desc;//description
    String showName;
    int type;

    public TextureInfo getIcon() {
        return icon;
    }

    public void setIcon(TextureInfo icon) {
        this.icon = icon;
    }

    int damageType;
    String buff;

    public int getDamageType() {
        return damageType;
    }

    public void setDamageType(int damageType) {
        this.damageType = damageType;
    }

    int castTime;
    int coolDown;
    TextureInfo icon;
    BoneBlock shape;
    int speed;
    int animation;
    public int getId() {
        return id;
    }

    public void setId(int id) {
        this.id = id;
    }

    public String getName() {
        return name;
    }

    public void setName(String name) {
        this.name = name;
    }

    public String getDesc() {
        return desc;
    }

    public void setDesc(String desc) {
        this.desc = desc;
    }

    public String getShowName() {
        return showName;
    }

    public void setShowName(String showName) {
        this.showName = showName;
    }

    public int getType() {
        return type;
    }

    public void setType(int type) {
        this.type = type;
    }

    public String getBuff() {
        return buff;
    }

    public void setBuff(String buff) {
        this.buff = buff;
    }

    public int getCastTime() {
        return castTime;
    }

    public void setCastTime(int castTime) {
        this.castTime = castTime;
    }

    public int getCoolDown() {
        return coolDown;
    }

    public void setCoolDown(int coolDown) {
        this.coolDown = coolDown;
    }

    public BoneBlock getShape() {
        return shape;
    }

    public void setShape(BoneBlock shape) {
        this.shape = shape;
    }

    public int getSpeed() {
        return speed;
    }

    public void setSpeed(int speed) {
        this.speed = speed;
    }

    public int getAnimation() {
        return animation;
    }

    public void setAnimation(int animation) {
        this.animation = animation;
    }

}


text:
<?php

return [
    'form_header' => [
        'email'       => 'Email change',
        'password'    => 'Reset Password',
        'profile'     => 'Your Profile',
        'information' => 'İnformation',
        'activitys'   => 'Activitys',
    ],
    'field_information' => [
        'email'       => 'Email of the account , which is provide to login in ',
        'password'    => 'Password of the account , which is provide to login in ',
        'profile'     => 'information about your profile',
        'information' => 'Your personal information',
    ],
    'field_name'=> [
        'password'   => 'Password',
        'name'       => 'Name',
        'about'      => 'About / Biografiy',
        'email'      => 'E-mail address',
        'cnfpassword'=> 'Confirm Password',
    ],
    'sidebar' => [
        'account'     => 'Account',
        'information' => 'İnformation',
        'activity'    => 'Activity',
        'logout'      => 'Logout',
    ],
    'table' => [
        'id'             => 'ID',
        'action'         => 'Action',
        'resources_name' => 'Resources name',
        'resources_id'   => 'Resources id',
        'action_time'    => 'Action Time',
    ],
    'button'=> [
        'update'=> 'update',
    ],
];


text:
/*

  This is a small tool to dump the contents of the input file.

  The code in this example file is released into the Public Domain.

*/

#include <iostream>

#include <osmium/handler/dump.hpp>
#include <osmium/io/any_input.hpp>

int main(int argc, char* argv[]) {
    std::ios_base::sync_with_stdio(false);

    if (argc < 2 || argc > 3) {
        std::cerr << "Usage: " << argv[0] << " OSMFILE [TYPES]\n";
        std::cerr << "TYPES can be any combination of 'n', 'w', 'r', and 'c' to indicate what types of OSM entities you want (default: all).\n";
        exit(1);
    }

    osmium::osm_entity_bits::type read_types = osmium::osm_entity_bits::all;

    if (argc == 3) {
        read_types = osmium::osm_entity_bits::nothing;
        std::string types = argv[2];
        if (types.find('n') != std::string::npos) read_types |= osmium::osm_entity_bits::node;
        if (types.find('w') != std::string::npos) read_types |= osmium::osm_entity_bits::way;
        if (types.find('r') != std::string::npos) read_types |= osmium::osm_entity_bits::relation;
        if (types.find('c') != std::string::npos) read_types |= osmium::osm_entity_bits::changeset;
    }

    osmium::io::Reader reader(argv[1], read_types);
    osmium::io::Header header = reader.header();

    std::cout << "HEADER:\n  generator=" << header.get("generator") << "\n";

    for (auto& bbox : header.boxes()) {
        std::cout << "  bbox=" << bbox << "\n";
    }

    osmium::handler::Dump dump(std::cout);
    while (osmium::memory::Buffer buffer = reader.read()) {
        osmium::apply(buffer, dump);
    }

    reader.close();
}


text:
This is a tiny plugins with less than 100 lines codes to help you easier to deal with DOM's className when your project don't need library like jQuery.
For modern browser, we use the HTML5 classList api to deal with className, otherwise, for browsers like IE9- or Android 3-, we use our function to deal with them.
## Support ##

	IE6+ & all mordern browser & all mobile browser.

## Quick Start ##

```javascript
//We set classList object to window, so you can use it directly like:

var dom = document.querySelector('#yourDom');
var className = "yourClassName";

classList.addClass(dom,className);
classList.removeClass(dom,className);
classList.toggleClass(dom,className);
classList.hasClass(dom,className);
```

## Parameters ##

```javascript
dom 	  : the DOM element that you need to set class name
className : the className that you want to use,
			you can also set multiple classes like : "newClass newClass2 newClass3"
```

## Method ##

```javascript

classList.addClass    : Add one or more class to the DOM you input.

classList.removeClass : Remove one or more class to the DOM you input.

classList.toggleClass : If the DOM you input has the className, remove it, otherwise, add the className.

classList.hasClass    : If the DOM has the className you input, return true, otherwise ,return false.
						(For this method, you can only input one className at one time.)

```

## More ##

1. You can set multiple classes at one time like :

```javascript
//add 3 classes at one time
classList.addClass(dom,"newClass newClass2 newClass3");
```
2. We support method chaining,so you can do something like this:

```javascript

classList.addClass(dom,'newClass').removeClass(dom,'oldClass').toggle(dom,'newClass oldClass');

```
If you find any issues, please contact me!
Hope this will help you.
Thanks!


text:
module GpasHelper
end


text:
package severity

import (
	"github.com/mediaFORGE/gol"
	"github.com/mediaFORGE/gol/fields/severity"
)

// Severity the struct for the filter.
type Severity struct {
	minimum severity.Type
}

// New creates a severity filter.
func New(s severity.Type) gol.LogFilter {
	return &Severity{
		minimum: s,
	}
}

// Filter performs a filter check on the given message.
// Returns whether or not a given message should be filtered.
func (f Severity) Filter(msg *gol.LogMessage) bool {

	if s, err := msg.Severity(); err != nil {
		// no severity
		return true
	} else {
		return s > f.minimum
	}
}

var _ gol.LogFilter = (*Severity)(nil)