# Using cl_float3 in parallel reduction example

Question asked by shunyo on May 17, 2013
Latest reply on May 20, 2013 by shunyo

I adapted the parallel reduction example for openCL for a bunch of floats. Now I wanted to expand the code to include cl_float3. So I want to find the minimum among a array of cl_float3. I thought it was a straight forward expansion from float to float3 in kernel. But I am receiving garbage values when i return from the kernel. Below is the kernel:

__kernel void pmin3(__global float3  *src,

__global float3  *gmin,

__local  float3  *lmin,

__global float  *dbg,

uint           nitems,

uint           dev)

{

uint count  = nitems     / get_global_size(0);

uint idx    = (dev == 0) ? get_global_id(0) * count

: get_global_id(0);

uint stride = (dev == 0) ? 1 : get_global_size(0);

// Private min for the work-item

float3 pmin = (float3)(pow(2.0,32.0)-1,pow(2.0,32.0)-1,pow(2.0,32.0)-1);

for (int n = 0; n < count; n++, idx += stride) {

pmin.x = min(pmin.x,src[idx].x);

pmin.y = min(pmin.y,src[idx].y);

pmin.z = min(pmin.z,src[idx].z);

}

// Reduce values within the work-group into local memory

barrier(CLK_LOCAL_MEM_FENCE);

if (get_local_id(0) == 0)

lmin = (float3)(pow(2.0,32.0)-1,pow(2.0,32.0)-1,pow(2.0,32.0)-1);

for (int n = 0; n < get_local_size(0); n++) {

barrier(CLK_LOCAL_MEM_FENCE);

if (get_local_id(0) == n) {

lmin.x = min(lmin.x,pmin.x);

lmin.y = min(lmin.y,pmin.y);

lmin.z = min(lmin.z,pmin.z);

}

}

barrier(CLK_LOCAL_MEM_FENCE);

// Write to __global gmin which will contain the work-group minima

if (get_local_id(0) == 0)

gmin[get_group_id(0)] = lmin;

// Collect debug information

if (get_global_id(0) == 0) {

dbg = get_num_groups(0);

dbg = get_global_size(0);

dbg = count;

dbg = stride;

}

}

__kernel void min_reduce3( __global float3  *gmin)

{

for (int n = 0; n < get_global_size(0); n++) {

barrier(CLK_GLOBAL_MEM_FENCE);

if (get_global_id(0) == n) {

gmin.x = min(gmin.x,gmin[n].x);

gmin.y = min(gmin.y,gmin[n].y);

gmin.z = min(gmin.z,gmin[n].z);

}

}

barrier(CLK_GLOBAL_MEM_FENCE);

}

I think it is the problem with get_global_id(0) and get_global_size() which gives the entire size instead of the only the number of rows to be given. Any suggestions?
{"ft_lang_label":"__label__en","ft_lang_prob":0.6566704,"math_prob":0.99785566,"size":2818,"snap":"2019-13-2019-22","text_gpt3_token_len":900,"char_repetition_ratio":0.13361762,"word_repetition_ratio":0.036319613,"special_character_ratio":0.34208658,"punctuation_ratio":0.20650406,"nsfw_num_words":0,"has_unicode_error":false,"math_prob_llama3":0.9891822,"pos_list":[0],"im_url_duplicate_count":[null],"WARC_HEADER":"{\"WARC-Type\":\"response\",\"WARC-Date\":\"2019-03-22T10:09:00Z\",\"WARC-Record-ID\":\"<urn:uuid:ee9ba234-48f5-4ae0-86e2-ca769aa12876>\",\"Content-Length\":\"146742\",\"Content-Type\":\"application/http; msgtype=response\",\"WARC-Warcinfo-ID\":\"<urn:uuid:6348b23c-e744-4d0e-9540-cff1a4ec9fce>\",\"WARC-Concurrent-To\":\"<urn:uuid:5b6b019d-2e08-4d23-94c1-38bbaf3909ea>\",\"WARC-IP-Address\":\"23.218.145.81\",\"WARC-Target-URI\":\"https://community.amd.com/thread/166666\",\"WARC-Payload-Digest\":\"sha1:3RR3ZXPAPFQZIS6DS7NCIOWWNE3Y2UQQ\",\"WARC-Block-Digest\":\"sha1:ODQAGL7LBOY7YWJYEEHKBIK2EMMAUHPF\",\"WARC-Identified-Payload-Type\":\"text/html\",\"warc_filename\":\"/cc_download/warc_2019/CC-MAIN-2019-13/CC-MAIN-2019-13_segments_1552912202642.32_warc_CC-MAIN-20190322094932-20190322120932-00049.warc.gz\"}"}