diff --git a/cl_kernel/kernel_hdr_lab.cl b/cl_kernel/kernel_hdr_lab.cl
new file mode 100644
index 0000000..39cf463
--- /dev/null
+++ b/cl_kernel/kernel_hdr_lab.cl
@@ -0,0 +1,58 @@
+/*
+ * function: kernel_hdr_lab
+ *     sample code of default kernel arguments
+ * input:    image2d_t as read only
+ * output:   image2d_t as write only
+ */
+"__constant float table[] = {4.3287616,5.9319224,7.1324205,8.1288157,8.9965763,9.7739191,10.483327,11.139331,11.751958,12.328467,12.874310,13.393700,13.889977,14.365837,14.823494,15.264792,15.691289,16.104307,16.504990,16.894327,17.273184,17.642323,18.002419,18.354071,18.697817,19.034143,19.363485,19.686239,20.002764,20.313389,20.618416,20.918123,21.212763,21.502573,21.787767,22.068552,22.345116,22.617630,22.886259,23.151157,23.412468,23.670324,23.924854,24.176174,24.424398,24.669630,24.911972,25.151518,25.388355,25.622572,25.854246,26.083458,26.310276,26.534771,26.757010,26.977057,27.194969,27.410807,27.624624,27.836473,28.046406,28.254467,28.254467,28.254467,28.254467,28.254467,28.254467,28.254467,28.254467,28.254467,28.377895,28.567524,28.755781,28.942692,29.128284,29.312582,29.495617,29.677410,29.857986,30.037369,30.215582,30.392645,30.568581,30.743410,30.917152,31.089828,31.261454,31.432049,31.601633,31.770222,31.937832,32.104481,32.270180,32.434952,32.598808,32.761761,32.923832,33.085026,33.245361,33.404850,33.563507,33.721340,33.878368,34.034599,34.190044,34.344715,34.498627,34.651783,34.804199,34.955887,35.106853,35.257107,35.406662,35.555523,35.703705,35.851208,35.998051,36.144238,36.289776,36.434673,36.578941,36.722588,36.865616,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.008038,37.013512,37.148350,37.282703,37.416573,37.549969,37.682888,37.815342,37.947330,38.078865,38.209946,38.340580,38.470768,38.600517,38.729836,38.858719,38.987175,39.115215,39.242832,39.370041,39.496838,39.623226,39.749214,39.874802,40,40.124805,40.249222,40.373260,40.496914,40.620193,40.743095,40.865631,40.987804,41.109608,41.231056,41.352146,41.472885,41.593269,41.713306,41.833000,41.952354,42.071369,42.190048,42.308392,42.426407,42.544094,42.661457,42.778500,42.895222,43.011627,43.127716,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.243496,43.315910,43.427536,43.538902,43.650013,43.760872,43.871479,43.981831,44.091938,44.201797,44.311413,44.420780,44.529911,44.638798,44.747448,44.855862,44.964039,45.071983,45.179695,45.287178,45.394428,45.501453,45.608253,45.714825,45.821178,45.927307,46.033215,46.138905,46.244377,46.349632,46.454674,46.559505,46.664120,46.768528,46.872723,46.976711,47.080494,47.184067,47.287441,47.390610,47.493576,47.596344,47.698910,47.801281,47.903454,48.005428,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.107212,48.116840,48.216503,48.316002,48.415333,48.514507,48.613514,48.712364,48.811050,48.909580,49.007950,49.106163,49.204220,49.302116,49.399860,49.497452,49.594887,49.692173,49.789303,49.886284,49.983109,50.079788,50.176319,50.272701,50.368935,50.465023,50.560966,50.656761,50.752411,50.847919,50.943279,51.038502,51.133583,51.228519,51.323318,51.417973,51.512493,51.606873,51.701115,51.795219,51.889191,51.983021,52.076717,52.170280,52.263706,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.357002,52.370274,52.462727,52.555069,52.647293,52.739407,52.831406,52.923294,53.015072,53.106739,53.198296,53.289738,53.381077,53.472301,53.563419,53.654430,53.745327,53.836124,53.926811,54.017391,54.107864,54.198231,54.288494,54.378651,54.468708,54.558655,54.648502,54.738243,54.827885,54.917419,55.006851,55.096188,55.185417,55.274551,55.363579,55.452511,55.541340,55.630070,55.718704,55.807240,55.895676,55.984013,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.072254,56.136337,56.224670,56.312920,56.401089,56.489174,56.577179,56.665100,56.752941,56.840698,56.928371,57.015968,57.103481,57.190918,57.278271,57.365547,57.452744,57.539856,57.626896,57.713852,57.800732,57.887531,57.974255,58.060902,58.147465,58.233959,58.320370,58.406708,58.492970,58.579155,58.665260,58.751289,58.837250,58.923130,59.008938,59.094666,59.180325,59.265907,59.351414,59.436852,59.522209,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.607498,59.676125,59.762394,59.848598,59.934742,60.020828,60.106846,60.192806,60.278702,60.364540,60.450314,60.536030,60.621685,60.707275,60.792812,60.878284,60.963699,61.049049,61.134342,61.219578,61.304752,61.389870,61.474926,61.559925,61.644867,61.729748,61.814568,61.899334,61.984039,62.068687,62.153282,62.237812,62.322292,62.406708,62.491070,62.575375,62.659622,62.743816,62.827950,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.912029,62.940441,63.026379,63.112267,63.198112,63.283909,63.369659,63.455360,63.541019,63.626637,63.712200,63.797718,63.883190,63.968624,64.054001,64.139343,64.224632,64.309875,64.395081,64.480240,64.565346,64.650414,64.735435,64.820412,64.905342,64.990234,65.075073,65.159874,65.244629,65.329346,65.414009,65.498634,65.583214,65.667747,65.752243,65.836693,65.921097,66.005463,66.089783,66.174057,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.258286,66.291786,66.378784,66.465752,66.552689,66.639587,66.726456,66.813293,66.900085,66.986855,67.073586,67.160286,67.246948,67.333572,67.420174,67.506737,67.593269,67.679764,67.766228,67.852654,67.939056,68.025421,68.111755,68.198051,68.284325,68.370560,68.456757,68.542923,68.629059,68.715164,68.801239,68.887276,68.973282,69.059258,69.145203,69.231110,69.316986,69.402832,69.488655,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.574432,69.659309,69.748878,69.838425,69.927956,70.017456,70.106934,70.196388,70.285820,70.375237,70.464622,70.553986,70.643333,70.732658,70.821953,70.911224,71.000481,71.089706,71.178917,71.268105,71.357262,71.446404,71.535530,71.624619,71.713692,71.802750,71.891777,71.980789,72.069771,72.158737,72.247673,72.336594,72.425499,72.514374,72.603226,72.692062,72.780876,72.869659,72.958427,73.047173,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.135902,73.151604,73.245392,73.339172,73.432938,73.526695,73.620430,73.714165,73.807884,73.901588,73.995285,74.088966,74.182640,74.276299,74.369942,74.463577,74.557198,74.650818,74.744415,74.838005,74.931580,75.025139,75.118698,75.212242,75.305771,75.399284,75.492798,75.586296,75.679771,75.773247,75.866707,75.960159,76.053596,76.147018,76.240433,76.333839,76.427231,76.520615,76.613983,76.707336,76.800690,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.894020,76.900002,77,77.100006,77.200005,77.300003,77.400002,77.500000,77.600006,77.700005,77.800003,77.900002,78,78.099998,78.199997,78.300003,78.400002,78.500000,78.599998,78.699997,78.799995,78.899994,79,79.099998,79.199997,79.299995,79.400002,79.500000,79.599998,79.699997,79.799995,79.900002,80,80.099998,80.199997,80.299995,80.400002,80.500000,80.599998,80.699997,80.800003,80.900002,81,81.099998,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.199997,81.299294,81.408134,81.516991,81.625862,81.734749,81.843643,81.952560,82.061493,82.170433,82.279388,82.388359,82.497353,82.606354,82.715363,82.824394,82.933441,83.042496,83.151573,83.260658,83.369759,83.478874,83.588005,83.697151,83.806305,83.915474,84.024666,84.133865,84.243080,84.352310,84.461555,84.570816,84.680084,84.789368,84.898666,85.007988,85.117310,85.226654,85.336006,85.445374,85.554764,85.664162,85.773575,85.882996,85.992439,86.101898,86.211365,86.320839,86.430336,86.539848,86.649376,86.758911,86.868454,86.978027,87.087608,87.197197,87.306801,87.416420,87.526054,87.635704,87.745361,87.855034,87.964722,88.074432,88.184143,88.293877,88.403618,88.513374,88.623146,88.732933,88.842728,88.952538,89.062363,89.172203,89.282051,89.391914,89.501793,89.611694,89.721596,89.831512,89.941444,90.051399,90.161354,90.271324,90.381310,90.491310,90.601326,90.711349,90.821388,90.931442,91.041512,91.151596,91.261681,91.371788,91.481911,91.592049,91.702187,91.812347,91.922516,92.032707,92.142906,92.253120,92.363342,92.473579,92.583839,92.694099,92.804375,92.914665,93.024971,93.135292,93.245621,93.355965,93.466316,93.576683,93.687073,93.797470,93.907875,94.018295,94.128731,94.239182,94.349640,94.460106,94.570595,94.681099,94.791611,94.902130,95.012665,95.123215,95.233788,95.344360,95.454948,95.565552,95.676170,95.786797,95.897430,96.008087,96.118752,96.229431,96.340126,96.450829,96.561539,96.672272,96.783012,96.893768,97.004532,97.115311,97.226112,97.336914,97.447731,97.558563,97.669403,97.780266,97.891129,98.002007,98.112900,98.223816,98.334732,98.445656,98.556602,98.667557,98.778526,98.889511,99.000504,99.111504,99.222519,99.333557,99.444603,99.555656,99.666718,99.777809,99.888893,100};   "
+
+"static float fun(float in)            "
+"{                                     "
+"    if(in > 0.008856)                 "
+"        return pow(in,1.0/3);          "
+"    else                              "
+"        return (float)(7.787*in + 16.0/116);    "
+"}                                               "
+"__kernel void kernel_hdr_lab (__read_only image2d_t input, __write_only image2d_t output)        "
+"{                                                                                             "
+"    int x = get_global_id (0);                                                                "
+"    int y = get_global_id (1);                                                                "
+"    sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;  "
+"    float4 pixel_in,pixel_out;									"
+"    unsigned int table_id;"
+"    pixel_in = read_imagef(input, sampler, (int2)(x,y));                                          "
+"    float X,Y,Z,L,a,b;                                                                                          "
+"    X = 0.412453*pixel_in.x + 0.357580*pixel_in.y + 0.180423*pixel_in.z;    "
+"    Y = 0.212671*pixel_in.x + 0.715160*pixel_in.y + 0.072169*pixel_in.z;    "
+"    Z = 0.019334*pixel_in.x + 0.119193*pixel_in.y + 0.950227*pixel_in.z;    "
+"    if(Y > 0.008856)             "
+"    L = 116.0*(pow(Y,1.0/3))  -16.0;    "
+"    else                       "
+"    L = 903.3*Y;               "
+"    a = 500*(fun(X) - fun(Y));    "
+"    b = 200*(fun(Y) - fun(Z));    "
+"    table_id = (unsigned int)(L*10) <1000 ? (unsigned int)(L*10) : 999;    "
+"    L = table[table_id];    "
+"    float fX,fY,fZ;    "
+"    Y = pow(L/116.0, 3.0);    "
+"    if (Y < 0.008856)    "
+"        Y = L/903.3;    "
+"    if (Y > 0.008856)    "
+"        fY = pow(Y, 1.0/3.0);    "
+"    else    "
+"        fY = 7.787*Y + 16.0/116.0;    "
+"    fX = a / 500.0 + fY;     "
+"    X = pow(fX, 3.0);    "
+"    if(X < 0.008865)    "
+"        X = (fX - 16.0/116.0) / 7.787;    "
+"    fZ = fY - b /200.0;    "
+"    Z = pow(fZ, 3.0);    "
+"    if(Z < 0.008865)    "
+"        Z = (fZ - 16.0/116.0) / 7.787;    "
+"    pixel_out.x = 3.240479*X - 1.537150*Y - 0.498535*Z;     "
+"    pixel_out.y = -0.969256*X + 1.875992*Y + 0.041556*Z;   "
+"    pixel_out.z = 0.055648*X - 0.204043*Y + 1.204043*Z;     "
+"    pixel_out.w = pixel_in.w;                             "
+"    write_imagef(output, (int2)(x,y), pixel_out);               "
+"}                                                                                             "
+
diff --git a/cl_kernel/kernel_hdr.cl b/cl_kernel/kernel_hdr_rgb.cl
similarity index 97%
rename from cl_kernel/kernel_hdr.cl
rename to cl_kernel/kernel_hdr_rgb.cl
index 86bd94c..3dde2f1 100644
--- a/cl_kernel/kernel_hdr.cl
+++ b/cl_kernel/kernel_hdr_rgb.cl
@@ -1,11 +1,11 @@
 /*
- * function: kernel_hdr
+ * function: kernel_hdr_rgb
  *     sample code of default kernel arguments
  * input:    image2d_t as read only
  * output:   image2d_t as write only
  */
 "	__constant int HDRTable[1024]={104,106,108,109,111,113,114,116,118,119,121,123,125,126,128,130,131,133,135,136,138,139,141,143,144,146,148,149,151,152,154,156,157,159,160,162,164,165,167,168,170,171,173,174,176,177,179,181,182,184,185,187,188,190,191,193,194,195,197,198,200,201,203,204,206,207,209,210,211,213,214,216,217,219,220,221,223,224,226,227,228,230,231,232,234,235,236,238,239,241,242,243,245,246,247,248,250,251,252,254,255,256,258,259,260,261,263,264,265,267,268,269,270,272,273,274,275,277,278,279,280,281,283,284,285,286,287,289,290,291,292,293,295,296,297,298,299,300,302,303,304,305,306,307,308,310,311,312,313,314,315,316,317,318,320,321,322,323,324,325,326,327,328,329,330,331,332,334,335,336,337,338,339,340,341,342,343,344,345,346,347,348,349,350,351,352,353,354,355,356,357,358,359,360,361,362,363,364,365,366,367,367,368,369,370,371,372,373,374,375,376,377,378,379,380,380,381,382,383,384,385,386,387,388,389,389,390,391,392,393,394,395,395,396,397,398,399,400,401,401,402,403,404,405,406,406,407,408,409,410,410,411,412,413,414,415,415,416,417,418,418,419,420,421,422,422,423,424,425,425,426,427,428,429,429,430,431,432,432,433,434,435,435,436,437,437,438,439,440,440,441,442,442,443,444,445,445,446,447,447,448,449,450,450,451,452,452,453,454,454,455,456,456,457,458,458,459,460,460,461,462,462,463,464,464,465,466,466,467,468,468,469,469,470,471,471,472,473,473,474,475,475,476,476,477,478,478,479,479,480,481,481,482,483,483,484,484,485,486,486,487,487,488,488,489,490,490,491,491,492,493,493,494,494,495,495,496,497,497,498,498,499,499,500,501,501,502,502,503,503,504,504,505,506,506,507,507,508,508,509,509,510,510,511,512,512,513,513,514,514,515,515,516,516,517,517,518,518,519,519,520,521,521,522,522,523,523,524,524,525,525,526,526,527,527,528,528,529,529,530,530,531,531,532,532,533,533,534,534,535,535,536,536,537,537,538,538,539,539,540,540,541,541,542,542,543,543,544,544,545,545,546,546,547,547,548,548,549,549,550,550,551,551,552,552,553,553,554,554,554,555,555,556,556,557,557,558,558,559,559,560,560,561,561,562,562,563,563,564,564,565,565,566,566,567,567,568,568,568,569,569,570,570,571,571,572,572,573,573,574,574,575,575,576,576,577,577,578,578,579,579,580,580,581,581,582,582,582,583,583,584,584,585,585,586,586,587,587,588,588,589,589,590,590,591,591,592,592,593,593,594,594,595,595,596,596,597,597,598,598,599,599,600,600,601,601,602,602,603,603,604,604,605,605,606,606,607,607,608,608,609,609,610,611,611,612,612,613,613,614,614,615,615,616,616,617,617,618,618,619,619,620,621,621,622,622,623,623,624,624,625,625,626,627,627,628,628,629,629,630,630,631,632,632,633,633,634,634,635,636,636,637,637,638,638,639,640,640,641,641,642,642,643,644,644,645,645,646,647,647,648,648,649,650,650,651,651,652,653,653,654,654,655,656,656,657,658,658,659,659,660,661,661,662,663,663,664,665,665,666,666,667,668,668,669,670,670,671,672,672,673,674,674,675,676,676,677,678,678,679,680,680,681,682,683,683,684,685,685,686,687,687,688,689,690,690,691,692,692,693,694,695,695,696,697,698,698,699,700,701,701,702,703,704,704,705,706,707,707,708,709,710,710,711,712,713,714,714,715,716,717,718,718,719,720,721,722,722,723,724,725,726,726,727,728,729,730,731,731,732,733,734,735,736,737,737,738,739,740,741,742,743,744,744,745,746,747,748,749,750,751,752,752,753,754,755,756,757,758,759,760,761,762,763,764,765,765,766,767,768,769,770,771,772,773,774,775,776,777,778,779,780,781,782,783,784,785,786,787,788,789,790,791,792,793,794,795,796,797,798,800,801,802,803,804,805,806,807,808,809,810,811,812,814,815,816,817,818,819,820,821,822,824,825,826,827,828,829,830,832,833,834,835,836,837,839,840,841,842,843,845,846,847,848,849,851,852,853,854,855,857,858,859,860,862,863,864,865,867,868,869,870,872,873,874,876,877,878,880,881,882,883,885,886,887,889,890,891,893,894,895,897,898,900,901,902,904,905,906,908,909,911,912,913,915,916,918,919,920,922,923,925,926,928,929,931,932,933,935,936,938,939,941,942,944,945,947,948,950,951,953,954,956,957,959,961,962,964,965,967,968,970,972,973,975,976,978,979,981,983,984,986,988,989,991,992,994,996,997,999,1001,1002,1004,1006,1007,1009,1011,1012,1014,1016,1017,1019,1021,1023};		"
-"__kernel void kernel_hdr (__read_only image2d_t input, __write_only image2d_t output)        "
+"__kernel void kernel_hdr_rgb (__read_only image2d_t input, __write_only image2d_t output)        "
 "{                                                                                             "
 "    int x = get_global_id (0);                                                                "
 "    int y = get_global_id (1);                                                                "
diff --git a/tests/test-cl-image.cpp b/tests/test-cl-image.cpp
index d3cb772..a306d0f 100644
--- a/tests/test-cl-image.cpp
+++ b/tests/test-cl-image.cpp
@@ -93,9 +93,8 @@
 static XCamReturn
 kernel_loop(SmartPtr<CLImageHandler> &image_handler, SmartPtr<DrmBoBuffer> &input_buf, SmartPtr<DrmBoBuffer> &output_buf, uint32_t kernel_loop_count)
 {
-    int i;
     XCamReturn ret = XCAM_RETURN_NO_ERROR;
-    for (i = 0; i < kernel_loop_count; i++) {
+    for (uint32_t i = 0; i < kernel_loop_count; i++) {
         PROFILING_START(cl_kernel);
         ret = image_handler->execute (input_buf, output_buf);
         PROFILING_END(cl_kernel, kernel_loop_count)
@@ -114,8 +113,10 @@
             "\t -i input     specify input file path\n"
             "\t -o output    specify output file path\n"
             "\t -p count     specify cl kernel loop count\n"
-            "\t -c csc_type    specify csc type\n"
-            "\t 			 select from [rgba2nv12, rgba2lab]\n"
+            "\t -c csc_type  specify csc type, default:rgba2nv12\n"
+            "\t              select from [rgba2nv12, rgba2lab]\n"
+            "\t -d hdr_type  specify hdr type, default:rgb\n"
+            "\t              select from [rgb, lab]\n"
             "\t -h           help\n"
             , bin_name);
 }
@@ -137,9 +138,10 @@
     SmartPtr<DrmDisplay> display;
     SmartPtr<DrmBoBufferPool> buf_pool;
     int opt = 0;
-    CLCscType csc_type = CL_CSC_TYPE_NONE;
+    CLCscType csc_type = CL_CSC_TYPE_RGBATONV12;
+    CLHdrType hdr_type = CL_HDR_TYPE_RGB;
 
-    while ((opt =  getopt(argc, argv, "f:i:o:t:p:c:h")) != -1) {
+    while ((opt =  getopt(argc, argv, "f:i:o:t:p:c:d:h")) != -1) {
         switch (opt) {
         case 'i':
             input_file = optarg;
@@ -155,6 +157,7 @@
                 format = V4L2_PIX_FMT_SGRBG10;
             else if (! strcasecmp (optarg, "rgba"))
                 format = V4L2_PIX_FMT_RGB32;
+
             else
                 print_help (bin_name);
             break;
@@ -189,6 +192,15 @@
             else
                 print_help (bin_name);
             break;
+        case 'd':
+            if (!strcasecmp (optarg, "rgb"))
+                hdr_type = CL_HDR_TYPE_RGB;
+            else if (!strcasecmp (optarg, "lab"))
+                hdr_type = CL_HDR_TYPE_LAB;
+            else
+                print_help (bin_name);
+            break;
+
         case 'h':
             print_help (bin_name);
             return 0;
@@ -235,7 +247,7 @@
         break;
     }
     case TestHandlerHDR:
-        image_handler = create_cl_hdr_image_handler (context);
+        image_handler = create_cl_hdr_image_handler (context, hdr_type);
         break;
     case TestHandlerWhiteBalance: {
         XCam3aResultWhiteBalance wb;
diff --git a/xcore/cl_3a_image_processor.cpp b/xcore/cl_3a_image_processor.cpp
index 17c9252..805a1eb 100644
--- a/xcore/cl_3a_image_processor.cpp
+++ b/xcore/cl_3a_image_processor.cpp
@@ -108,7 +108,7 @@
 
     /* hdr */
     if (_enable_hdr) {
-        image_handler = create_cl_hdr_image_handler (context);
+        image_handler = create_cl_hdr_image_handler (context, CL_HDR_TYPE_RGB);
         _hdr = image_handler;
         XCAM_FAIL_RETURN (
             WARNING,
diff --git a/xcore/cl_hdr_handler.cpp b/xcore/cl_hdr_handler.cpp
index fc0cdaa..ddbcbfd 100644
--- a/xcore/cl_hdr_handler.cpp
+++ b/xcore/cl_hdr_handler.cpp
@@ -22,85 +22,34 @@
 
 namespace XCam {
 
-CLHdrImageKernel::CLHdrImageKernel (SmartPtr<CLContext> &context)
-    : CLImageKernel (context, "kernel_hdr")
-{
-}
-
-XCamReturn
-CLHdrImageKernel::prepare_arguments (
-    SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output,
-    CLArgument args[], uint32_t &arg_count,
-    CLWorkSize &work_size)
-{
-    SmartPtr<CLContext> context = get_context ();
-    const VideoBufferInfo & video_info = input->get_video_info ();
-    cl_libva_image image_info;
-    uint32_t channel_bits = XCAM_ALIGN_UP (video_info.color_bits, 8);
-
-    xcam_mem_clear (&image_info);
-    image_info.fmt.image_channel_order = CL_R;
-    if (channel_bits == 8)
-        image_info.fmt.image_channel_data_type = CL_UNORM_INT8;
-    else if (channel_bits == 16)
-        image_info.fmt.image_channel_data_type = CL_UNORM_INT16;
-    image_info.offset = 0;
-    image_info.width = video_info.width;
-    image_info.height = (video_info.size / video_info.strides[0])/4*4;
-    image_info.row_pitch = video_info.strides[0];
-
-    _image_in = new CLVaImage (context, input, &image_info);
-    _image_out = new CLVaImage (context, output, &image_info);
-
-    XCAM_ASSERT (_image_in->is_valid () && _image_out->is_valid ());
-    XCAM_FAIL_RETURN (
-        WARNING,
-        _image_in->is_valid () && _image_out->is_valid (),
-        XCAM_RETURN_ERROR_MEM,
-        "cl image kernel(%s) in/out memory not available", get_kernel_name ());
-
-    //set args;
-    args[0].arg_adress = &_image_in->get_mem_id ();
-    args[0].arg_size = sizeof (cl_mem);
-    args[1].arg_adress = &_image_out->get_mem_id ();
-    args[1].arg_size = sizeof (cl_mem);
-    arg_count = 2;
-
-    work_size.dim = XCAM_DEFAULT_IMAGE_DIM;
-    work_size.global[0] = image_info.row_pitch;
-    work_size.global[1] = image_info.height;
-    work_size.local[0] = 8;
-    work_size.local[1] = 4;
-
-    return XCAM_RETURN_NO_ERROR;
-}
-
-
-XCamReturn
-CLHdrImageKernel::post_execute ()
-{
-    return CLImageKernel::post_execute ();
-}
-
 SmartPtr<CLImageHandler>
-create_cl_hdr_image_handler (SmartPtr<CLContext> &context)
+create_cl_hdr_image_handler (SmartPtr<CLContext> &context, CLHdrType type)
 {
     SmartPtr<CLImageHandler> hdr_handler;
     SmartPtr<CLImageKernel> hdr_kernel;
     XCamReturn ret = XCAM_RETURN_NO_ERROR;
 
-    hdr_kernel = new CLHdrImageKernel (context);
-    {
-        XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_hdr)
-#include "kernel_hdr.cl"
-        XCAM_CL_KERNEL_FUNC_END;
-        ret = hdr_kernel->load_from_source (kernel_hdr_body, strlen (kernel_hdr_body));
-        XCAM_FAIL_RETURN (
-            WARNING,
-            ret == XCAM_RETURN_NO_ERROR,
-            NULL,
-            "CL image handler(%s) load source failed", hdr_kernel->get_kernel_name());
+    XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_hdr_rgb)
+#include "kernel_hdr_rgb.cl"
+    XCAM_CL_KERNEL_FUNC_END;
+    XCAM_CL_KERNEL_FUNC_SOURCE_BEGIN(kernel_hdr_lab)
+#include "kernel_hdr_lab.cl"
+    XCAM_CL_KERNEL_FUNC_END;
+
+    if (type == CL_HDR_TYPE_RGB) {
+        hdr_kernel = new CLImageKernel (context, "kernel_hdr_rgb");
+        ret = hdr_kernel->load_from_source (kernel_hdr_rgb_body, strlen (kernel_hdr_rgb_body));
     }
+    else if (type == CL_HDR_TYPE_LAB) {
+        hdr_kernel = new CLImageKernel (context, "kernel_hdr_lab");
+        ret = hdr_kernel->load_from_source (kernel_hdr_lab_body, strlen (kernel_hdr_lab_body));
+    }
+    XCAM_FAIL_RETURN (
+        WARNING,
+        ret == XCAM_RETURN_NO_ERROR,
+        NULL,
+        "CL image handler(%s) load source failed", hdr_kernel->get_kernel_name());
+
     XCAM_ASSERT (hdr_kernel->is_valid ());
     hdr_handler = new CLImageHandler ("cl_handler_hdr");
     hdr_handler->add_kernel  (hdr_kernel);
diff --git a/xcore/cl_hdr_handler.h b/xcore/cl_hdr_handler.h
index 85bba38..06fd026 100644
--- a/xcore/cl_hdr_handler.h
+++ b/xcore/cl_hdr_handler.h
@@ -26,25 +26,13 @@
 
 namespace XCam {
 
-class CLHdrImageKernel
-    : public CLImageKernel
-{
-public:
-    explicit CLHdrImageKernel (SmartPtr<CLContext> &context);
-
-    virtual XCamReturn post_execute ();
-protected:
-    virtual XCamReturn prepare_arguments (
-        SmartPtr<DrmBoBuffer> &input, SmartPtr<DrmBoBuffer> &output,
-        CLArgument args[], uint32_t &arg_count,
-        CLWorkSize &work_size);
-
-private:
-    XCAM_DEAD_COPY (CLHdrImageKernel);
+enum CLHdrType {
+    CL_HDR_TYPE_RGB,
+    CL_HDR_TYPE_LAB,
 };
 
 SmartPtr<CLImageHandler>
-create_cl_hdr_image_handler (SmartPtr<CLContext> &context);
+create_cl_hdr_image_handler (SmartPtr<CLContext> &context, CLHdrType type);
 
 };