cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

edwen
Journeyman III

how to initialize a constant variable of array?

constant memory initialization

I am trying to modify a CUDA program into OpenCL and I met a problem. In my "market.cl" file, I need to declare and initialize some constant variables. for example, "__constant int v1 = 5;" This works fine. However, I need to declare some arrays, such as "__constant float delta[5][80]". In my CUDA program, I used the function "cudaMemcpyToSymbol" to copy all constants into constant memory, and the code in the "market.cu" is something like:

#define NumFactors 8

float delta_[5][80];

float weight = sqrt(1.0f/float)NumFactors);

for (int i = 0; i < 5; i++)

{  for (int j = 0; j < 80; j++)

         delta_ = weight;

}

cudaMemcpyToSymbol (delta, &delta_, sizeof(delta_);

 

How can I realize this in OpenCL? Thank you very much for your advice.

 

0 Likes
11 Replies
genaganna
Journeyman III

Originally posted by: edwen I am trying to modify a CUDA program into OpenCL and I met a problem. In my "market.cl" file, I need to declare and initialize some constant variables. for example, "__constant int v1 = 5;" This works fine. However, I need to declare some arrays, such as "__constant float delta[5][80]". In my CUDA program, I used the function "cudaMemcpyToSymbol" to copy all constants into constant memory, and the code in the "market.cu" is something like:

#define NumFactors 8

float delta_[5][80];

float weight = sqrt(1.0f/float)NumFactors);

for (int i = 0; i < 5; i++)

{  for (int j = 0; j < 80; j++)

         delta_ = weight;

}

cudaMemcpyToSymbol (delta, &delta_, sizeof(delta_);

How can I realize this in OpenCL? Thank you very much for your advice.

 



Setting __constant variable array from runtime is exactly same as global memory variable.

1. Create a buffer and assign values from delta_ array.

2. Set kernel argument to appropriate kernel parameter.

0 Likes

Setting __constant variable from running is exactly same as global memory.

1. Create a buffer and assign values from delta_ array.

2. Set kernel argument to appropriate kernel parameter.

Thanks for your prompt reply. However, in my "market.cl", besides two kernel functions, there are several other functions. I need to use the "delta" array in these functions, and then the two kernel functions call these functions.  Therefore, I can not just pass the array as argument to certain kernel parameter. The structure of my "market.cl" is somehow like:

float func1(...)

{ ... //code using array "delta"}

float func2(...)

{.../code using array "delta"}

__kernel func (...)

{...;

y1 = func1(...);

y2 = func2(...)}

there would be a run-time error if I pass array "delta" to the kernel function "func" as argument since "func1" and "func2" need this array.

 

0 Likes

Originally posted by: edwen


 

Setting __constant variable from running is exactly same as global memory.

 

1. Create a buffer and assign values from delta_ array.

 

2. Set kernel argument to appropriate kernel parameter.

 

 

Thanks for your prompt reply. However, in my "market.cl", besides two kernel functions, there are several other functions. I need to use the "delta" array in these functions, and then the two kernel functions call these functions.  Therefore, I can not just pass the array as argument to certain kernel parameter. The structure of my "market.cl" is somehow like:

 

float func1(...)

 

{ ... //code using array "delta"}

 

float func2(...)

 

{.../code using array "delta"}

 

__kernel func (...)

 

{...;

 

y1 = func1(...);

 

y2 = func2(...)}

 

there would be a run-time error if I pass array "delta" to the kernel function "func" as argument since "func1" and "func2" need this array.

 

 

Just initialize your delta_ array in .cl file.  Initializing __constant array is just simple to general C-array

__constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

 

0 Likes

Originally posted by: genaganna
Originally posted by: edwen


Just initialize your delta_ array in .cl file.  Initializing __constant array is just simple to general C-array

__constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

 

I thought about this way. However, too reasons this doesn't work for me that well. Firstly, the array is very big, and I need to use a for loop to generate the values. Secondly, in the for loop generating the values of the array element, some math functions, such as sqrt and exp, are used. It's not applicable to initialize the array using your method. I wonder if there is an alternative way in OpenCL. Thanks.

0 Likes

Originally posted by: edwen
Originally posted by: genaganna
Originally posted by: edwen


 

Just initialize your delta_ array in .cl file.  Initializing __constant array is just simple to general C-array

 

__constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

 

 

 

 

I thought about this way. However, too reasons this doesn't work for me that well. Firstly, the array is very big, and I need to use a for loop to generate the values. Secondly, in the for loop generating the values of the array element, some math functions, such as sqrt and exp, are used. It's not applicable to initialize the array using your method. I wonder if there is an alternative way in OpenCL. Thanks.

 

1. Initialize values of delta_ in host and create CLBuffer and assign this values to CLBuffer

2. All kernels should have a __constant parameter so that you can pass your value to all kernels.

 

what is size of your __constant buffer size.

0 Likes

Originally posted by: genaganna
Originally posted by: edwen
Originally posted by: genaganna
Originally posted by: edwen


1. Initialize values of delta_ in host and create CLBuffer and assign this values to CLBuffer

2. All kernels should have a __constant parameter so that you can pass your value to all kernels.

 

what is size of your __constant buffer size.

The problem is, in my .cl file, besides the kernel functions, there are some other functions using this array. If I only pass the array to the kernel functions as argument, those non-kernel functions will not be able to use the array.

0 Likes

Originally posted by: edwen

 

The problem is, in my .cl file, besides the kernel functions, there are some other functions using this array. If I only pass the array to the kernel functions as argument, those non-kernel functions will not be able to use the array.

 

I said all kernels should have this as parameter.

0 Likes

Originally posted by: genaganna
Originally posted by: edwen

 

The problem is, in my .cl file, besides the kernel functions, there are some other functions using this array. If I only pass the array to the kernel functions as argument, those non-kernel functions will not be able to use the array.

 

I said all kernels should have this as parameter.

ok, I think i got your point. So I can pass the array to the kernel functions frist. And when the kernel functions call the non-kernel functions, it pass the array to the non-kernel functions. Is that right? Remember, we can not pass any arguments to the non-kernel functions in the .cl directly from the host.

0 Likes

Originally posted by: edwen

 

ok, I think i got your point. So I can pass the array to the kernel functions frist. And when the kernel functions call the non-kernel functions, it pass the array to the non-kernel functions. Is that right?

Yes you are right.

Remember, we can not pass any arguments to the non-kernel functions in the .cl directly from the host.

 

Yes, it is not possible to pass any arguments to non-kernel functions.

0 Likes

Originally posted by: genaganna
Originally posted by: edwen

 

ok, I think i got your point. So I can pass the array to the kernel functions frist. And when the kernel functions call the non-kernel functions, it pass the array to the non-kernel functions. Is that right?

Yes you are right.

Remember, we can not pass any arguments to the non-kernel functions in the .cl directly from the host.

 

Yes, it is not possible to pass any arguments to non-kernel functions.

This method is so inefficient. I have more than 10 constant variables in my .cl file, which are used by those non-kernal functions. Now I have to add corresonding parameters to all these functions. And accordingly, in my kernel functions, add extra parameters while calling these non-kernel functions. I think the CUDA method is much better, namely we only declare the constant variables in .cl file, but calculate the values of these variables in the host, and then copy them into the constant memory through sth similar to cudaMemcpyToSymbol(). OpenCL should improve the way it declares and initializes variables in constant memories.

0 Likes

There is one other possibility, in case you are using run-time compilation from source in your opencl code. You can generate the code for initializing the constant varible in the main program, add it to the kernel code, and then compile it. Once you have all your data for the constants, you can for sure write a function which generates this kind of code as a string:

__constant float delta_[8][10] = {{2, 2, ..}, {2, 2, ..} ..}

Then prefix your opencl code with this string and compile.

OpenCL is run-time compiled - use that for your best advantage! Think more about that... you can replace parameters with compile-time defines (allowing the compiler to unroll loops), etc.

0 Likes