cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

edwen
Journeyman III

Can non-kernel functions access global memory?

My .cl is somehow like this:

void evaluation(float *L)

{...}

__kernel void memdata(__global float *d_v, __global float LC, const int num)

{ const int tid = get_global_id(0);

__global float *L_loc = LC + num*tid;

evaluation(L_loc);

 }

 

This always brings me a programbuild error. The reason I have to add "__global" in front of L_loc is because LC is in global memory, otherwise L_loc can't point to LC at all. However, "evaluation" is a non-kernel function, and it seems it can't reach global memory. Is that right?

0 Likes
9 Replies
n0thing
Journeyman III

This compiles fine after changing the argument pointer of the function to point to global memory.

void evaluation(__global float *L)

{}

__kernel void memdata(__global float *d_v, __global float *LC, const int num)

{ const int tid = get_global_id(0);

__global float *L_loc = LC + num*tid;

evaluation(L_loc);

 }

0 Likes

Yes, I tried that solution, too. It compiles fine, but fails when I try to run it.

0 Likes

Fails as in correctness failure? Can you post the complete code?

BTW The following code works just fine for me -

unsigned int func(__global unsigned int *ptr, unsigned int mul) { return ptr[get_global_id(0)] * mul; } __kernel void templateKernel(__global unsigned int * output, __global unsigned int * input, const unsigned int multiplier) { uint tid = get_global_id(0); output[tid] = func(input,multiplier); }

0 Likes

Originally posted by: n0thing Fails as in correctness failure? Can you post the complete code?

BTW The following code works just fine for me -

 

Thanks again. Let me try your code and get back to you.

0 Likes

Originally posted by: n0thing This compiles fine after changing the argument pointer of the function to point to global memory.

void evaluation(__global float *L)

Seems like in a non-kernel function, it's not allow to add __global in front a parameter. You can try to add any irrelevant function with that kind of parameter in you own code and see. Thanks,

 

0 Likes

I can confirm I'm using __global parameters in non-kernel functions successfully. So maybe try out changing your code to something following: 

void evaluation(__global float *L, const int startIndex) { //acces the array as L[startIndex] } __kernel void memdata(__global float *d_v, __global float *LC, const int num) { const int tid = get_global_id(0); evaluation(LC, num*tid); }

0 Likes

Thanks, n0thing and karbous. Yes, n0thing's code works fine. I will figure out what's wrong with mine. I couldn't post the whole thing before i get the permission from the project manager.

0 Likes

Originally posted by: karbous I can confirm I'm using __global parameters in non-kernel functions successfully. So maybe try out changing your code to something following: 

I think I figured out what my problem is. Please look at the code:


float valuation(const int t, float *L, const int numFactors_, const int numLibors_, const int numTimeSteps_, const InstrumentType instrumentType_)
{float *npv;
float npv_;
npv_ = 0.0f;
npv = &npv_;

if (instrumentType_ == SWAP)
{swap_valuation(t, L, npv);}

else if (instrumentType_ == CAP)
{cap_valuation(t, L, npv);}

else if (instrumentType_ == SWAPTION)
{swaption_valuation(t, L, npv, numFactors_, numLibors_, numTimeSteps_);
}
return *npv;
}

__kernel void fun1(__global float *d_v,...)
{...
evaluation(...);
}

__kernel void fun2(__global float *d_v, __global float *L,...)
{...
evaluation(...);
}

As in above, I have two kernel functions, and they both call the "evalution" function. However, in kernel one, I dont have a global array argument named "L" passed from my host code, and the "L" array required by function "evaluation" will be generated locally. In the second kernel function, I do have an global argument named "L" passed from the host code. If I add "__global" to the parameter "L" in function "evaluation", the first kernel function wont work any more. Further more, all the functions inside "evaluation" would have problems, too, since they dont have a "__global" sign, either.

OpenCL should really allows overriding.

 



0 Likes

You can create a__global array inside the kernel 1 instead of __local and then there would be no pointer conversion issues.

Or have two differrent functions.

Thanks.

0 Likes