fix compile errors on Linux (#130)

* fix compile errors on Linux

* remove warnings

Author: luoyu-intel

include/experimental/group/gemm/impl/int4_dequantize_xmx_xe.hpp

@@ -530,7 +530,7 @@ class gemm_t<compute_policy_int4_dequantize_xmx<compute_attr_,
                     xetla_vector<uint8_t, block_size_x_b * block_size_y_b>
                             zero_pt_blk;
 #pragma unroll
-                    for (int row = 0; row < block_size_y_b; row++) {
+                    for (uint32_t row = 0; row < block_size_y_b; row++) {
                         zero_pt_blk
                                 .xetla_select<block_size_x_b, 1>(
                                         row * block_size_x_b)

include/subgroup/tile/impl/prefetch_xe.hpp

@@ -152,7 +152,7 @@ tile_prefetch(payload_t &payload) {
             = payload_t::arch_tag == gpu_arch::Xe ? 64 : 32;
     if constexpr (prefetch_len >= reg_in_bytes) {
 #pragma unroll
-        for (int j = 0; j < prefetch_len / reg_in_bytes; j++) {
+        for (uint32_t j = 0; j < prefetch_len / reg_in_bytes; j++) {
             uint32_t offset_x = j * reg_in_bytes * payload_t::scale_factor;
             uint32_t address_offset = offset_x * sizeof(dtype);
             xetla_prefetch_global<prefetch_dtype, reg_in_bytes,

tests/integration/data_transformer/main.cpp

@@ -86,10 +86,10 @@ static void data_transformer_run() {
     auto inputBundle
             = get_kernel_bundle<bundle_state::input>(context, kernelId);
     static const std::string env_set_str = "SYCL_PROGRAM_COMPILE_OPTIONS= -vc-codegen -Xfinalizer ' -printregusage -enableBCR -DPASTokenReduction '";
-    putenv(env_set_str.c_str());
+    putenv(const_cast<char*>(env_set_str.c_str()));
     kernel_bundle<bundle_state::executable> exeBundle = build(inputBundle);
     static const std::string env_unset_str = "SYCL_PROGRAM_COMPILE_OPTIONS=";
-    putenv(env_unset_str.c_str());
+    putenv(const_cast<char*>(env_unset_str.c_str()));
 
     try {
         auto e_esimd = queue.submit([&](handler &cgh) {

tests/integration/gemm/bf16_stream_k/main.cpp

@@ -301,7 +301,7 @@ void stream_k_gemm_run(uint32_t iter) {
 
 
     static const std::string env_set_str = "SYCL_PROGRAM_COMPILE_OPTIONS= -vc-codegen -doubleGRF -vc-disable-indvars-opt -Xfinalizer ' -printregusage -enableBCR -DPASTokenReduction '";
-    putenv(env_set_str.c_str());
+    putenv(const_cast<char*>(env_set_str.c_str()));
     //Define and initialize the data required for the calculation
     auto A = alloc_device_and_init<data_type_a>(
             size_a,
@@ -434,7 +434,7 @@ void stream_k_gemm_run(uint32_t iter) {
     }
 
     static const std::string env_unset_str = "SYCL_PROGRAM_COMPILE_OPTIONS=";
-    putenv(env_unset_str.c_str());
+    putenv(const_cast<char*>(env_unset_str.c_str()));
 
     ASSERT_EQ(0,
             gemm_result_validate(A, B, C, Bias, matrix_m, matrix_k, matrix_n,

tests/integration/softmax/softmax_bwd.cpp

@@ -97,11 +97,11 @@ void softmax_bwd_run() {
                 context, kernelId);
 
         static const std::string env_set_str = "SYCL_PROGRAM_COMPILE_OPTIONS= -vc-codegen -doubleGRF  -Xfinalizer ' -printregusage -enableBCR -DPASTokenReduction '";
-        putenv(env_set_str.c_str());
+        putenv(const_cast<char*>(env_set_str.c_str()));
         sycl::kernel_bundle<sycl::bundle_state::executable> exeBundle
                 = build(inputBundle);
         static const std::string env_del_str = "SYCL_PROGRAM_COMPILE_OPTIONS=";
-        putenv(env_del_str.c_str());
+        putenv(const_cast<char*>(env_del_str.c_str()));
         try {
 
             auto e_softmax_bwd = queue.submit([&](sycl::handler &cgh) {

tests/integration/softmax/softmax_fwd.cpp

@@ -83,10 +83,10 @@ void softmax_fwd_run() {
         auto inputBundle
                 = get_kernel_bundle<bundle_state::input>(context, kernelId);
         static const std::string env_set_str = "SYCL_PROGRAM_COMPILE_OPTIONS= -vc-codegen -doubleGRF -Xfinalizer ' -printregusage -enableBCR -DPASTokenReduction '";
-        putenv(env_set_str.c_str());
+        putenv(const_cast<char*>(env_set_str.c_str()));
         kernel_bundle<bundle_state::executable> exeBundle = build(inputBundle);
         static const std::string env_del_str = "SYCL_PROGRAM_COMPILE_OPTIONS=";
-        putenv(env_del_str.c_str());
+        putenv(const_cast<char*>(env_del_str.c_str()));
         try {
 
             auto e_softmax_fwd = queue.submit([&](handler &cgh) {

tests/unit/tile_load_store/kernel_func.hpp

@@ -128,12 +128,12 @@ struct tile_load_store_unaligned_2d_func {
         payload_store_t payload_store(mem_desc_c);
         tile_prefetch(payload_prefetch);
         tile_load(matA, payload_load);
-        static const char vec_a[] = "after load %f %f %f %f\n";
+        // static const char vec_a[] = "after load %f %f %f %f\n";
         // sycl::ext::oneapi::experimental::printf(vec_a, (matA.reg[0]), (matA.reg[1]), (matA.reg[2]), (matA.reg[3]));
         // sycl::ext::oneapi::experimental::printf(vec_a, (matA.reg[16]), (matA.reg[17]), (matA.reg[18]), (matA.reg[19]));
         matC.reg = matA.reg;
         tile_store(matC, payload_store);
-        static const char vec_b[] = "%d %d %d %d\n";
+        // static const char vec_b[] = "%d %d %d %d\n";
         // sycl::ext::oneapi::experimental::printf(vec_b, (matC.reg[0]), (matC.reg[1]), (matC.reg[2]), (matC.reg[3]));
     }
 };

tests/utils/execution.hpp

@@ -90,10 +90,10 @@ void gemm_exec(const std::string &compile_str, size_t batch = 1) {
         auto inputBundle
                 = get_kernel_bundle<bundle_state::input>(context, kernelId);
         static const std::string env_set_str = "SYCL_PROGRAM_COMPILE_OPTIONS="+compile_str;
-        putenv(env_set_str.c_str());
+        putenv(const_cast<char*>(env_set_str.c_str()));
         kernel_bundle<bundle_state::executable> exeBundle = build(inputBundle);
         static const std::string env_unset_str = "SYCL_PROGRAM_COMPILE_OPTIONS=";
-        putenv(env_unset_str.c_str());
+        putenv(const_cast<char*>(env_unset_str.c_str()));
 
         using namespace gpu::xetla::group;
         using namespace gpu::xetla::kernel;